1
0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-26 08:28:20 +00:00
Commit Graph

1995 Commits

Author SHA1 Message Date
jsteube
81a447b167 Fixed a function declaration attribute in -m 8900 kernel leading to unuseable -m 9300 which shares kernel code with -m 8900 2018-06-21 13:46:53 +02:00
jsteube
32d6b3e10e OpenCL kernels: Add '-pure' prefix to kernel filenames to avoid problems caused by reusing existing hashcat installation folder 2018-06-20 14:18:17 +02:00
jsteube
547025ec47 HCCAPX management: Use advanced hints in message_pair stored by hcxtools about endian bitness of replay counter
Fixed missing code section in -m 2500 and -m 2501 to crack corrupted handshakes with a LE endian bitness base
2018-06-15 17:00:41 +02:00
Mathieu Geli
4dbc1f4a87 Implement 7701/7801 SAP CODVN half-hashes 2018-03-06 16:42:53 +03:00
jsteube
8079abffb0 Fixed a missing kernel in -m 5600 in combination with -a 3 and -O if mask is >= 16 characters 2018-02-28 11:25:52 +01:00
Jens Steube
a71c69983d Make words_buf_r in DES bitsliced kernels __constant 2018-02-21 10:50:24 +01:00
jsteube
ca1115a1ee No longer need to use 32 threads on second dimension for bitsliced algorithms 2018-02-20 01:01:50 +01:00
Jens Steube
ad50883080 Allow unroll for DES based algorithms but not bitsliced versions 2018-02-18 11:28:25 +01:00
Jens Steube
e79feb0b6f Add more reqd_work_group_size attributes to kernels 2018-02-17 22:16:05 +01:00
Jens Steube
ea2f158cf8 Give JiT a hint about bcrypt running at 8 threads always 2018-02-17 21:33:11 +01:00
Jens Steube
81b229c08a Make new c_append_helper a bit more AMD friendly 2018-02-17 15:18:19 +01:00
jsteube
3a23b275e5 Improved c_append_helper[] handling 2018-02-17 14:24:29 +01:00
jsteube
64eb9ca9ef Fix md5crypt speed on GTX1080 2018-02-17 11:45:05 +01:00
Jens Steube
aa82d8d34d Re-enable CPU optimizations and some CPU case in thread management 2018-02-16 18:56:21 +01:00
Jens Steube
483ae613ee Copy/paste error in optimized -m 2500 AUX 3 kernel 2018-02-16 11:48:11 +01:00
Jens Steube
a15c165773 Revert the pos == 0 change due to some unexpected results on nvidia 2018-02-16 11:38:08 +01:00
jsteube
874635cc49 Do not use a vector function to write into a scalar variable even if vectorize support is disabled 2018-02-16 10:48:49 +01:00
jsteube
71adf1bd09 Do not use a vector function to write into a scalar variable even if vectorize support is disabled 2018-02-16 10:28:16 +01:00
jsteube
f596d076aa Optimize some WPA comparison kernel 2018-02-16 00:30:47 +01:00
Jens Steube
e0e796bc2d More optimized -m 500, -m 1600 and -m 6300 pure kernel 2018-02-15 15:35:22 +01:00
Jens Steube
97f569a8ca More optimized -m 500, -m 1600 and -m 6300 pure kernel 2018-02-15 12:31:59 +01:00
Jens Steube
f6f16f56af A bit optimized -m 500 pure kernel 2018-02-14 19:13:23 +01:00
jsteube
b626e7f61b Optimize update functions in inc_hash_* helper files by testing for pos == 0 2018-02-14 17:31:41 +01:00
Jens Steube
ec63c2f017 switch_buffer_* function can return sooner if offset is zero 2018-02-14 11:35:00 +01:00
Jens Steube
2dddef839c Fix sha1_update_64 debugging comment 2018-02-14 11:03:35 +01:00
Jens Steube
2a19f19904 Drop -m 6800 from no-unroll list in inc_vendor.cl 2018-02-13 17:08:29 +01:00
jsteube
dfb95024bc Fix temporary datatype in wpapmk kernels 2018-02-13 16:34:00 +01:00
jsteube
fe4413797e OpenCL Kernels: Use three separate comparison kernels (depending on keyver) for WPA instead of one 2018-02-13 09:13:35 +01:00
jsteube
00bd356ade Synchronize m01000s with m00900s 2018-02-12 09:03:11 +01:00
jsteube
aa65ed28e5 Switch back c_append_helper to static 2018-02-12 09:02:37 +01:00
Jens Steube
d656e9c3a4 OpenCL Kernels: Use the kernel local buffer size as additional reference in order to limit the thread-count 2018-02-11 10:56:08 +01:00
jsteube
5951207365 Get rid of some old volatiles 2018-02-09 19:18:30 +01:00
jsteube
05a01d3843 fix some datatypes 2018-02-08 19:13:29 +01:00
jsteube
d5153539e2 Some syntax error 2018-02-08 09:49:59 +01:00
jsteube
786384664e DECLSPEC for CPU 2018-02-08 09:42:59 +01:00
jsteube
4cbd0eb812 Fix missing compressor kernel in --stdout mode 2018-02-07 22:28:52 +01:00
jsteube
512fb5f6fb No inline keyword for rules 2018-02-07 15:02:58 +01:00
jsteube
e4e1c1d515 We can't mix inline functions with static constants 2018-02-07 14:16:27 +01:00
Jens Steube
8273bb8376 NV JiT doesn't like static inline keywords 2018-02-06 22:05:15 +01:00
jsteube
3e08750900 OpenCL Kernels: Add general function declaration keyword (static inline) 2018-02-06 19:12:24 +01:00
jsteube
5391edca0d Weird macOS JiT likes this more 2018-02-05 19:39:20 +01:00
jsteube
a3a16f676f OpenCL Kernels: Add a decompressing kernel and a compressing host code in order to reduce PCIe transfer time
For details see https://hashcat.net/forum/thread-7267.html
2018-02-05 17:18:58 +01:00
jsteube
53e2b40bad Fixed a uninitialized value in OpenCL kernels 9720, 9820 and 10420 leading to absurd benchmark performance 2018-02-02 14:02:33 +01:00
jsteube
13a79cf942 Fix unoptimized Kerberos 5 TGS-REP etype 23 kernel for use on macOS 2018-01-31 21:28:46 +01:00
jsteube
e877c30ebc OpenCL Kernels: Remove password length restriction to 16 for Cisco-PIX and Cisco-ASA hashes
Fixes #1488
2018-01-27 22:21:44 +01:00
jsteube
18bb0a9493 Add missing gpu_atinit() to m02000_a0 and m02000_a1 2018-01-26 20:19:25 +01:00
jsteube
6a04e953e0 Fix missing gpu_atinit in --stdout kernel 2018-01-26 14:56:57 +01:00
jsteube
a9d5f571b7 Remove extra token at end of #include directive 2018-01-25 23:48:31 +01:00
Jens Steube
1f1eacca95 Fix -m 16600 -a 3 optimize mode kernels function declaration for CPU use 2018-01-25 19:04:30 +01:00
jsteube
553668bb9f Added hash-mode 16600 = Electrum Wallet (Salt-Type 1-3) 2018-01-25 15:28:21 +01:00
Jens Steube
7062425d2b OpenCL Kernels: Use a special kernel to initialize the password buffer used during autotune measurements to reduce startup time 2018-01-23 20:33:26 +01:00
jsteube
0796c074c3 Added -m 16500 Kernels
Also changed function declaration of parser function from const hashconfig_t to just hashconfig_t
2018-01-21 18:53:55 +01:00
jsteube
ee9ec0f9a7 Add JWT esalt datatype 2018-01-21 15:32:37 +01:00
Jens Steube
bb806d777e
Merge pull request #1493 from mohemiv/master
added -m 16400 = CRAM-MD5 Dovecot
2018-01-19 10:13:46 +01:00
jsteube
53f3da9f63 OpenCL Kernels: Use static declaraction for uXXa variables used in __constant space 2018-01-18 23:19:31 +01:00
jsteube
3a303ffce4 Replace variables from uXX to uXXa if used in __constant space 2018-01-18 23:16:11 +01:00
Arseniy Sharoglazov
928cf471fb The hash-mode for "CRAM-MD5 Dovecot" changed from 10201 to 16400 2018-01-17 11:25:21 +03:00
Arseniy Sharoglazov
798f05355f added -m 10201 = CRAM-MD5 Dovecot 2018-01-15 15:52:52 +03:00
philsmd
bf656774bb
fixes #1279: added -m 16300 = Ethereum Pre-Sale Wallet, PBKDF2-HMAC-SHA256 2017-12-20 11:41:46 +01:00
jsteube
0d89ddfcd9 Finish adding hash-mode 16200 = Apple Secure Notes 2017-12-13 12:32:38 +01:00
jsteube
f573c1d96d Add optimized -m 16100 kernels 2017-12-03 14:35:39 +01:00
jsteube
9a3cf88887 Update -m 16100 kernel to enable cracking of sequences 1, 3 and 5 2017-11-30 13:41:25 +01:00
jsteube
e5ca2e2fcb Add more kernels for -m 16100 2017-11-30 10:16:14 +01:00
jsteube
5847067c96 First working -m 16100 kernel 2017-11-29 17:00:14 +01:00
jsteube
1b312d14fd Added hash-mode 16000 = Tripcode 2017-11-11 14:44:56 +01:00
jsteube
00abb849e1 Update inc_vector.cl for ideal performance with reference GTX1080 for NV and Vega64 for AMD 2017-10-26 13:41:47 +02:00
jsteube
503f676cb0 Remove static keywords from OpenCL functions, they can cause old NV JiT compiler to fail 2017-10-20 13:23:43 +02:00
jsteube
9c832092df Fix some function declarations inside LUKS kernel 2017-10-10 01:27:32 +02:00
jsteube
b4c3df876c Fix some formating in -m 600 2017-10-09 18:45:11 +02:00
jsteube
7c2dadab17 no volatile required here 2017-09-21 20:08:24 +02:00
Jens Steube
55eaff3a45 Merge pull request #1365 from Fist0urs/DPAPI_reworked
-m 15300 reworked + splitted in 15300/15900
2017-09-21 17:02:18 +02:00
Fist0urs
e3cb3e9b4c test.pl ready and 0 error. Ready for PR 2017-09-21 16:55:30 +02:00
jsteube
7a17b8159b all() function is not working as expected in scalar datatype case 2017-09-21 12:23:33 +02:00
Fist0urs
a6294537fd Splitted DPAPI kernel in 2 to increase performances 2017-09-21 12:23:33 +02:00
jsteube
207ce9b853 all() function is not working as expected in scalar datatype case 2017-09-20 23:00:00 +02:00
jsteube
ddbe805c00 Fix last step of make_kn 2017-09-20 09:40:14 +02:00
jsteube
68f5b12754 Get rid of swap32() in make_kn() in -m 2500 2017-09-19 13:45:43 +02:00
jsteube
ab1dabebbe Fix missing include in -m 2501 2017-09-19 12:19:46 +02:00
jsteube
b14f44dcf7 Fix uninitialized keymic buffer 2017-09-19 11:58:18 +02:00
jsteube
beab5457e6 Backport WPA-PSK-SHA256-AES-CMAC to -m 2501 2017-09-19 10:22:03 +02:00
jsteube
ca1b6492e7 Some code simplify on AES CMAC 2017-09-19 01:12:29 +02:00
jsteube
4e3a642f7f Initial WPA2-PSK-SHA256-AES-CMAC support 2017-09-19 01:08:38 +02:00
jsteube
98fc02e04b Add PTK compute for keyver 3 2017-09-18 13:50:09 +02:00
jsteube
617dbb97ba Prepare migration -m 15800 into -m 2500 2017-09-18 13:21:00 +02:00
philsmd
2dadae4e9a fixed incorrect use of the esalt_bufs for -m 600 = BLAKE2-512 2017-09-17 15:28:24 +02:00
mhasbini
de7ccd88ef Fix overflow in mangle_dupechar_last function 2017-09-16 20:43:38 +03:00
jsteube
b169653b8f Fix missing return value in rule_op_mangle_toggle_at() 2017-09-08 22:49:49 +02:00
jsteube
55f653f374 Get rid of volatile in TrueCrypt kernels 2017-09-08 19:47:56 +02:00
jsteube
16e33b20fc Fix out of boundary access in -m 4700 2017-09-08 19:42:34 +02:00
jsteube
51dd982b12 Bring back some volatile for AMD 2017-09-08 14:08:21 +02:00
jsteube
9125062ffc Move volatiles for AMD closer to the problem 2017-09-08 13:32:19 +02:00
jsteube
1963b12fdc According to AMD docs, GCN 3 and 4 are the same 2017-09-08 12:46:00 +02:00
jsteube
ac9f1da747 Add fine-tuned AMD GCN control macros 2017-09-07 20:33:43 +02:00
Jens Steube
a910aea9e0 Do not use __local memory for -m 13100 if running on a device without physical shared memory 2017-09-05 17:41:55 +02:00
Jens Steube
b58aa445b4 Do not use __local memory for -m 7500 if running on a device without physical shared memory 2017-09-05 17:37:20 +02:00
Jens Steube
63f6ca5114 Do not use __local memory for whirlpool if running on a device without physical shared memory 2017-09-05 16:45:20 +02:00
jsteube
8b0e7087c7 Fixed an invalid optimization code in kernel 7700 depending on the input hash, causing the kernel to loop forever 2017-09-03 13:36:14 +02:00
jsteube
151dbc5349 Fix replace value in inc_hash_ripemd160.cl 2017-09-01 16:35:08 +02:00
jsteube
f859f466ef Fix -m 8300 in -a 0 mode 2017-09-01 16:10:29 +02:00
jsteube
f5e04254dc Fix -m 10800 in -a 0 mode 2017-09-01 16:06:42 +02:00
jsteube
d3b9febb30 Fix some double variable declarations 2017-08-30 16:29:25 +02:00
jsteube
40b57677cd OpenCL Kernels: Reactivate Dalibors XOR optimization on MD5_H on all MD5 based algorithms 2017-08-30 15:32:09 +02:00
jsteube
6d112aeb39 OpenCL Kernels: Rewritten Keccak kernel to run fully on registers and partially reversed last round 2017-08-30 13:27:04 +02:00
jsteube
a378abee66 Add missing NEW_SIMD_CODE in -m 6600 2017-08-29 12:01:43 +02:00
jsteube
1c169af0ad Make -m 14100 a pure kernel only 2017-08-28 22:26:30 +02:00
jsteube
2b9888486e Make -m 14000 a pure kernel only and add volatile for asm statement 2017-08-28 22:20:40 +02:00
jsteube
99f416435e Fix invalid use of __constant in LM kernel 2017-08-28 19:40:51 +02:00
jsteube
6db2f4cc18 Fix typo 2017-08-28 15:54:47 +02:00
jsteube
918578bee1 Improve some NVidia specific inline assembly 2017-08-28 14:15:47 +02:00
jsteube
9de1e557bb More VEGA specific inline assembly to improve SHA1 based kernels 2017-08-28 09:24:06 +02:00
jsteube
a0be36d7b8 Fix compile error caused by __add3() 2017-08-27 19:46:17 +02:00
jsteube
00e38cc2c6 Add VEGA specific inline assembly to improve all MD4, MD5, SHA1 and SHA256 based kernels 2017-08-27 19:36:07 +02:00
jsteube
7bfd343ec9 Optimized rule_op_mangle_dupechar_last(), rule_op_mangle_rotate_right(), rule_op_mangle_rotate_left() and append_block1() in rule engine 2017-08-27 16:47:21 +02:00
jsteube
52a97fee75 Improve rule engine performance by improving append_0x80_xxx() performance by using precomputed values from constant memory 2017-08-27 14:22:20 +02:00
jsteube
3260000357 Fix whirlpool pure kernel in -a 0 mode 2017-08-26 19:51:37 +02:00
jsteube
e3810d054b Fix some use of pw_t tmp variable 2017-08-26 19:48:38 +02:00
jsteube
5e01ff4c53 Refactor some u32x to u32 where u32x is not needed 2017-08-26 18:31:50 +02:00
jsteube
1aa76eac15 Refactor use of __constant to match up with the user selected attack mode 2017-08-25 17:52:55 +02:00
jsteube
938c281ee0 Resurrect some volatile variables in order to correctly compile pure kernels on AMD drivers 2017-08-25 17:06:07 +02:00
jsteube
48fbe81a09 Add more inline assembly for AMD ROCm 2017-08-25 16:33:00 +02:00
jsteube
6c619155c3 Workaround ROCm compiler error in aes256_ExpandKey() 2017-08-25 12:10:36 +02:00
jsteube
8c9c36ee2a Fix out-of-bound access in aesXXX_InvertKey() 2017-08-25 11:52:07 +02:00
jsteube
bed7e8f466 Remove unused truncate_block_xxx_xx() functions and update kernels to use the _S function 2017-08-24 20:07:43 +02:00
jsteube
51dc1c7db3 Use truncate_block_4x4_le_S() instead of truncate_block_4x4_le() in -m 6800 2017-08-24 19:53:29 +02:00
jsteube
9b73c464d2 Fix typo in macro 2017-08-24 17:19:16 +02:00
jsteube
7b443ee7ff Optimize performance of rule_op_mangle_title_sep(), rule_op_mangle_purgechar() and rule_op_mangle_replace() 2017-08-24 17:14:33 +02:00
jsteube
0de41c2716 Some more optimizations for rule engine 2017-08-24 15:09:55 +02:00
jsteube
9f8c5a253d More rule engine performance optimizations 2017-08-24 00:49:46 +02:00
jsteube
0783289e2f Optimized a0 pure kernel for AMD 2017-08-23 13:40:22 +02:00
jsteube
a5659d5619 Also switch optimized kernels rule engine to make use of kernel rules in constant memory 2017-08-23 12:46:14 +02:00
jsteube
1d04de3a8e Limit kernel-loops in straight-mode to 256, therefore allow rules to be stored in constant memory 2017-08-23 12:43:59 +02:00
jsteube
51372438fe Allow OpenCL kernel inline assembly if ROCm drivers was detected 2017-08-22 18:47:53 +02:00
jsteube
8853884f2a Fix append_four_byte() in case sm8 is 0 2017-08-21 16:04:43 +02:00
jsteube
f32e113942 Add missing case in append_block() in pure kernel rule engine 2017-08-20 15:08:51 +02:00
jsteube
6907981f08 Backport current state of optimized kernel rule engine to CPU 2017-08-20 12:50:24 +02:00
jsteube
508f1562f2 Fix --stdout kernels, gid_max was still set to u32 2017-08-20 12:13:34 +02:00
jsteube
319799bbbf Switch the datatypes of the variables responsible for work-item count and work-item size from u32 to u64 2017-08-19 16:39:22 +02:00
jsteube
d9c906e134 Move 0x80 to hardcoded position for sha3-256 bit in order to allow ROCm compiler to use registers only 2017-08-18 16:22:25 +02:00
jsteube
694cc0b740 Remove all calls to overwrite_at_* functions 2017-08-17 16:20:01 +02:00
jsteube
e984a829ea Remove no longer needed overwrite_at_* functions 2017-08-17 15:53:09 +02:00
jsteube
bf299fe043 Optimized 3DES for rocm 2017-08-17 14:03:55 +02:00
jsteube
ad1ce462d1 Get rid of ceil() in OpenCL kernels 2017-08-17 13:43:35 +02:00
jsteube
53f53fe014 Reduced number of required registers in SIP based on maximum possible esalt length 2017-08-17 12:16:49 +02:00
jsteube
9ee5da40e0 Workaround rocm compiler error for -m 15300 2017-08-17 11:25:34 +02:00
jsteube
88e995ddcf Replace some SIMD related function calls 2017-08-17 11:18:39 +02:00
jsteube
5b5bdf3889 Replace some SIMD related function calls 2017-08-17 10:18:17 +02:00
jsteube
967e96728d Make all the OpenCL kernel function includes static 2017-08-16 20:27:17 +02:00
jsteube
21e9c63d46 Fix rotl64() the same was as rotr64() 2017-08-16 17:58:33 +02:00
jsteube
58012ada0c Fall back to old rotr64 optimization for AMD 2017-08-16 16:14:46 +02:00
philsmd
4a89172140
reformatting; replaced some tabs with spaces 2017-08-16 13:46:40 +02:00
jsteube
ec874c1d59 Optimized the following pure kernel rule engine functions:
- mangle_lrest()
- mangle_lrest_ufirst()
- mangle_urest()
- mangle_urest_lfirst()
- mangle_trest()
- mangle_toggle_at()
- mangle_reverse()
- mangle_dupeword()
- mangle_reflect()
- mangle_rotate_left()
- mangle_rotate_right()
- mangle_switch_first()
- mangle_switch_last()
- mangle_switch_at()
- mangle_title_sep()
- mangle_title_sep()
Added some helper functions:
- generate_cmask()
- append_four_byte()
- append_three_byte()
- append_two_byte()
- append_one_byte()
- append_block()
- exchange_byte()
Removed some helper functions:
- upper_at()
- lower_at()
- toggle_at()
- mangle_switch()
NOTE: Changes need to be backported to CPU when finished
2017-08-13 16:43:46 +02:00
jsteube
9a57c4b20e Fix missing boundary check in pure kernel rule engine 2017-08-12 13:32:05 +02:00
jsteube
dcaa91a88f Fix rule engine function call in amp_a0 2017-08-12 13:28:31 +02:00
jsteube
0b0abb5c12 Prepare pure kernel rule engine for performance optimization 2017-08-12 13:04:52 +02:00
jsteube
98b4aab9d0 Update inc_rp_optimized to inc_rp for pure kernels 2017-08-11 23:51:17 +02:00
jsteube
cb1fe251de Add missing functions to pure kernel rule engine 2017-08-11 22:58:37 +02:00
jsteube
07b54c1257 Replace code to use pure kernel rule engine for slow hashes 2017-08-11 16:21:19 +02:00
jsteube
8a0d21360b Prepare for pure kernel rule engine support 2017-08-11 16:09:12 +02:00
jsteube
34d882a116 Rename inc_rp.X to inc_rp_optimized.X 2017-08-11 11:25:47 +02:00
jsteube
68a8f70edb Mix in pure kernel functions in various optimized kernels 2017-08-10 13:56:53 +02:00
jsteube
66f7590883 Mix in pure kernel functions in various optimized kernels 2017-08-10 12:38:42 +02:00
jsteube
a228e296da Fix some whirlpool vector function declaration 2017-08-10 12:18:55 +02:00
jsteube
560a786ea7 Mix in pure kernel functions in various optimized kernels 2017-08-10 12:04:07 +02:00
jsteube
4b2f3011e9 Mix in pure kernel functions to PDF 1.7 Level 8 (Acrobat 10 - 11) 2017-08-10 10:43:30 +02:00
jsteube
a956a84edb Fix vector datatype in -m 2810 2017-08-10 09:56:02 +02:00
jsteube
e6cb69e4cb Add pure kernels for JKS Java Key Store Private Keys (SHA1) 2017-08-09 15:16:12 +02:00
jsteube
adacccecdf Add pure kernels for FileZilla Server >= 0.9.55 2017-08-09 14:48:47 +02:00
jsteube
4c7f61e473 Add pure kernels for sha1(CX) 2017-08-09 14:26:39 +02:00
jsteube
b4c9f46205 Prepare sha1(CX) optimized kernel for pure kernel version 2017-08-09 13:21:33 +02:00
jsteube
29e13d6b77 Add pure kernels for OpenCart 2017-08-09 13:17:15 +02:00
jsteube
b6cf3144de Prepare OpenCart optimized kernel for pure kernel version 2017-08-09 12:35:01 +02:00
jsteube
4443ecd861 Add pure kernels for Windows Phone 8+ PIN/password 2017-08-08 15:26:34 +02:00
jsteube
a5c0aa6041 Add pure kernels for PeopleSoft PS_TOKEN 2017-08-08 11:45:22 +02:00
jsteube
24a2fb01aa Fix missing barrier in -m 8500 2017-08-07 18:58:23 +02:00
jsteube
0a676b549f Remove global barrier when not needed to workaround Intel OpenCL runtime bug 2017-08-07 17:25:15 +02:00
jsteube
b9876c100b Add pure kernels for AxCrypt in-memory SHA1 2017-08-07 15:28:23 +02:00
jsteube
5c6b3fa7ab Add pure kernels for Kerberos 5 TGS-REP etype 23 2017-08-07 15:22:18 +02:00
jsteube
51128473bc Add pure kernels for ColdFusion 10+ 2017-08-07 14:22:15 +02:00
jsteube
4f72c8bee6 Add pure kernels for SIP digest authentication (MD5) 2017-08-07 13:39:17 +02:00
jsteube
b1f9ed4a7c Add pure kernels for MySQL CRAM (SHA1) 2017-08-06 15:33:38 +02:00
jsteube
7548e5f85a Add pure kernels for PostgreSQL CRAM (MD5) 2017-08-06 13:54:02 +02:00
jsteube
c5c12f89c1 Rewrite code around amd_bytealign to be of type BE to save a branch afterwards 2017-08-05 19:46:56 +02:00
jsteube
a53d9e09de Fix some issue with offset_minus_4 2017-08-04 14:12:58 +02:00
jsteube
c9cae1f663 Add pure kernels for PrestaShop 2017-08-04 13:03:54 +02:00
jsteube
bc9f721dcd Mix in pure kernel functions to PDF 1.1 - 1.3 (Acrobat 2 - 4) 2017-08-03 20:30:46 +02:00
jsteube
177800d1d0 Add pure kernels for RAdmin2 2017-08-03 15:21:39 +02:00
jsteube
3a042972a4 Mix in pure kernel functions to MS Office <= 2003 $3/$4, SHA1 + RC4 2017-08-03 14:52:09 +02:00
jsteube
a650b0864e Mix in pure kernel functions to MS Office <= 2003 $0/$1, MD5 + RC4 2017-08-03 14:44:09 +02:00
jsteube
37432b19bc Mix in pure kernel functions to MS Office <= 2003 $0/$1, MD5 + RC4 2017-08-03 14:42:11 +02:00
jsteube
14983a7542 Simplify RACF kernel 2017-08-03 14:27:53 +02:00
jsteube
1f42377931 Simplify Lotus Notes/Domino 5 kernel 2017-08-03 14:11:31 +02:00
jsteube
c68191e47a Add pure kernels for DNSSEC (NSEC3) 2017-08-03 13:05:30 +02:00
jsteube
54eb0b158d Prepare DNSSEC (NSEC3) optimized kernel for pure kernel version 2017-08-03 12:35:05 +02:00
jsteube
c2a8ae0207 Add pure kernels for WBB3 (Woltlab Burning Board) 2017-08-03 12:21:53 +02:00
jsteube
7cf3c29ef5 Mix in pure kernel functions to DNSSEC (NSEC3) 2017-08-02 14:34:36 +02:00
jsteube
aafda5fa1b Add pure kernels for Citrix NetScaler 2017-08-02 14:12:27 +02:00
jsteube
5da64a1a43 Mix in pure kernel functions to SAP CODVN F/G (PASSCODE) 2017-08-02 13:24:54 +02:00
jsteube
bc6b8ca1c9 Mix in pure kernel functions to SAP CODVN B (BCODE) 2017-08-02 13:14:01 +02:00
jsteube
89d52f8209 Add pure kernels for Kerberos 5 AS-REQ Pre-Auth etype 23 2017-08-01 17:39:12 +02:00
jsteube
af6052d34b Revert some invalid rename of kernel files 2017-08-01 15:01:16 +02:00
jsteube
2802f1d592 Fix vector function calls and datatypes 2017-08-01 14:56:09 +02:00
jsteube
1eb249c5b4 Add pure kernels for IPMI2 RAKP HMAC-SHA1 2017-08-01 14:42:28 +02:00
jsteube
83d37ebeff Add pure kernels for FortiGate (FortiOS) 2017-08-01 14:16:27 +02:00
jsteube
cbd8f81a1c Add pure kernels for RipeMD160 2017-08-01 09:59:20 +02:00
jsteube
6946329b02 Fix BF pure kernels for NetNTLMv2 2017-07-31 15:37:49 +02:00
jsteube
a9fed50ce0 Add pure kernels for NetNTLMv2 2017-07-31 15:29:28 +02:00
jsteube
443fa960d3 Fix vector function calls 2017-07-31 15:28:22 +02:00
jsteube
50aeade65c Add pure kernels for NetNTLMv1 / NetNTLMv1+ESS 2017-07-31 10:23:04 +02:00
jsteube
5d137ba036 Add pure kernels for IKE-PSK SHA1 2017-07-31 09:36:57 +02:00
jsteube
942b7068be Add pure kernels for IKE-PSK MD5 2017-07-29 23:19:15 +02:00
jsteube
b541e46b9b Add pure kernels for Half MD5 2017-07-29 18:53:08 +02:00
jsteube
e0a565234a Optimized -m 7700 for ROCm 2017-07-28 20:52:53 +02:00
jsteube
332396a003 Fix SCRYPT on ROCm 2017-07-28 02:28:52 +02:00
jsteube
a85be1d0f0 Fix some const keywords in inc_truecrypt_xts.cl 2017-07-24 14:46:58 +02:00
jsteube
02e2279d59 Optimized -m 8500 for ROCm 2017-07-24 14:33:34 +02:00
jsteube
5bcda7d05a Optimized -m 5300 and -m 5400 for ROCm 2017-07-24 13:18:38 +02:00
jsteube
772441448a Optimized -m 8000 for ROCm 2017-07-24 13:13:35 +02:00
jsteube
9562d07264 Replace bitwise swaps with rotate() versions for AMD 2017-07-23 17:01:15 +02:00
jsteube
3125a756d9 Remove some AMD _unroll restrictions no longer required with ROCm 2017-07-23 14:44:20 +02:00
jsteube
4dca908cdf Fix a typo in OpenCL/m01460_a3-optimized.cl 2017-07-23 14:06:32 +02:00
jsteube
4c71bc984e Fix const keywords in -m 8600 2017-07-23 13:01:54 +02:00
jsteube
c255a967df Fix some types in rotate functions 2017-07-22 18:59:01 +02:00
jsteube
c289fcb2b0 Merge branch 'master' of https://github.com/hashcat/hashcat 2017-07-22 18:06:41 +02:00
jsteube
5e34ec348e Optimize kernels for ROCm 1.6
- Remove inline keywords
- Remove volatile keywords where it causes ROCm to slow down
- Replace DES functions (looks like bitselect somehow is no longer mapped to BFI_INT)
2017-07-22 18:05:18 +02:00
jsteube
27edc07c2f Add pure kernels for iSCSI CHAP authentication, MD5(CHAP) 2017-07-21 16:59:10 +02:00
jsteube
e9821a01ba Add pure kernels for sha1($salt.$pass.$salt) 2017-07-21 16:42:55 +02:00
jsteube
9515927cf7 Add pure kernels for sha1(md5()) 2017-07-21 15:44:29 +02:00
jsteube
1fdb9d1d7e Add pure kernels for sha1($salt.sha1($pass)) 2017-07-20 18:06:54 +02:00
jsteube
15d725e6cc Add pure kernels for sha1(sha1($pass)) 2017-07-20 17:50:07 +02:00
jsteube
4e97a4db24 Add pure kernels for md5(sha1($pass)) 2017-07-20 17:38:43 +02:00
jsteube
03bb234045 Preparation for WPA/WPA2 AES-CMAC: works till PMK 2017-07-20 12:46:18 +02:00
jsteube
441434840c Fix broken -m 7900 after migration to pure kernel 2017-07-19 17:45:58 +02:00
jsteube
eae9329761 Workaround some AMD JiT compiler segfault on complex kernels 2017-07-19 13:34:36 +02:00
jsteube
920911b56e Migrate MD5-PIX and MD5-ASA to run as optimized kernels 2017-07-18 15:53:25 +02:00
jsteube
beb6ee2061 Add OPTI_TYPE_OPTIMIZED_KERNEL
Rename unconverted fast hash kernels to optimized kernels
Finalize some converted fast hashes to default kernels
2017-07-18 14:45:15 +02:00
jsteube
fbea72ebd6 Renamed default kernels to optimized kernels
Renamed pure kernels to default kernels
Replaced long option --length-limit-disable with --optimized-kernel-enable
Replaced short option -L with -O
Set --optimized-kernel-enable to unset by default
2017-07-18 13:23:42 +02:00
jsteube
2616d54794 Add -L kernel for md5(strtoupper(md5($pass))) 2017-07-18 12:01:18 +02:00
jsteube
0e7d40dd3d Add -L kernel for md5($salt.md5($pass.$salt)) 2017-07-18 11:52:14 +02:00
jsteube
ad9ee5c5b1 Add -L kernel for md5($salt.md5($salt.$pass)) 2017-07-18 11:35:38 +02:00
jsteube
80324918e9 Add -L kernel for md5(md5($pass).md5($salt)) 2017-07-17 20:18:56 +02:00
jsteube
945cf9be2f md5($salt.$pass.$salt) 2017-07-17 17:24:32 +02:00
jsteube
50760b57e4 Add -L kernel for md5(.md5()), MediaWiki B type 2017-07-17 15:27:37 +02:00
jsteube
d08c7689bb Add -L kernel for md5(md5(s).md5(p)), MyBB 1.2+, IPB2+ 2017-07-17 10:59:09 +02:00
jsteube
90a8f58459 Add -L kernel for md5(md5()), vBulletin, PHPS 2017-07-17 10:18:26 +02:00
jsteube
e863a12624 Some fixes for inc_hash_sha224.cl and inc_hash_sha384.cl 2017-07-17 08:40:36 +02:00
jsteube
fe38379d0d Add -L kernel for SHA224 2017-07-17 08:29:46 +02:00
jsteube
dbf74b68b2 Add -L kernel for Domain Cached Credentials (DCC), MS Cache 2017-07-16 19:38:38 +02:00
jsteube
cc4ff214d6 Add -L kernel for NTLM 2017-07-16 16:50:36 +02:00
jsteube
b2f3bfb06c Add -L kernel for MD4 2017-07-16 16:44:53 +02:00
Jens Steube
3d9b071e1e Improve CPU cracking speed by replacing vector comparison functions with the more advanced ones available on CPU 2017-07-16 12:30:09 +02:00
jsteube
2dd1833998 Move from ld.global.v4.u32 to ld.const.v4.u32 in _a3 kernels 2017-07-15 13:35:24 +02:00
jsteube
cd313c9c28 Add -L kernel for MySQL4.1/MySQL5 2017-07-14 23:58:47 +02:00
jsteube
757f3a39c2 Accidentially pushed experimental -m 2500 kernel 2017-07-14 23:10:05 +02:00
jsteube
8434e451ef Add -L support for all SHA512 based generic hashes 2017-07-14 17:25:23 +02:00
jsteube
7205f450dd Backport more HMAC functions in inc_hash_xxx.cl from global to private 2017-07-14 16:58:30 +02:00
jsteube
0334cdb277 Fix some datatypes 2017-07-14 16:44:30 +02:00
jsteube
5c75eb84b0 Add -L support for all SHA256 based generic hashes 2017-07-14 15:02:21 +02:00
jsteube
e2371540e0 Add missing kernel in -m m00150_a1-pure.cl and m00160_a1-pure.cl 2017-07-14 14:26:37 +02:00
jsteube
c4098e2230 Fix invalid use of a non-vector function from within a vector function 2017-07-14 14:16:48 +02:00
jsteube
e4683aebb8 Add -L support for all sha1 based generic hashes 2017-07-14 13:40:14 +02:00
jsteube
696afc2a1b Fix selector in switch_buffer_by_offset_1x64_be_S() 2017-07-14 13:37:23 +02:00
jsteube
4e0972ce3a Add xxx_update_vector_swap(), xxx_update_vector_utf16le_swap() and xxx_update_vector_utf16beN() for later use 2017-07-14 13:24:40 +02:00
jsteube
9c6c21490f Add *_hmac_init_swap for later use 2017-07-13 19:22:31 +02:00
jsteube
f03156b05e Add switch_buffer_by_offset_1x64_be_S() and code generators for later use 2017-07-13 18:46:24 +02:00
jsteube
5707fed499 Add example -L kernel for algorithms using HMAC where the password is the data 2017-07-13 15:33:25 +02:00
jsteube
994e7efc91 Add example -L kernel for algorithms using HMAC where the password is the key 2017-07-13 15:03:26 +02:00
jsteube
9c12459852 Add HMAC vector functions to inc_hash_* 2017-07-13 12:18:17 +02:00
jsteube
8537ae9f8e Add example -L kernel for algorithms with prepended salt in utf16le 2017-07-13 11:24:54 +02:00
jsteube
2191af7a8e Fix datatype in sha384_update_vector_utf16le() sha512_update_vector_utf16le() 2017-07-13 11:07:56 +02:00
jsteube
c512e0c01a Add example -L kernel for algorithms with appended salt in utf16le 2017-07-13 00:16:29 +02:00
jsteube
c082bea018 Add missing swaps to -m 6100 2017-07-12 23:23:04 +02:00
jsteube
5437de75ea Add some missing swaps to -m 1700 and -m 10800 2017-07-12 16:36:38 +02:00
jsteube
ec816485c9 Add example -L kernel for algorithms with appended salt 2017-07-12 16:15:37 +02:00
jsteube
0113aedfdd Unify some variable names in pure kernels 2017-07-12 15:58:58 +02:00
jsteube
facf7ad534 Add example -L kernel for algorithms with prepended salt 2017-07-12 15:52:31 +02:00
jsteube
9b6c6df53d Add xxx_nit_vector_from_scalar() to all inc_hash_xxx.cl includes 2017-07-12 15:45:22 +02:00
jsteube
728be2f587 Add missing -a0 and -a1 -L kernel for -m 6100 and -m 10800 2017-07-12 14:50:48 +02:00
jsteube
4b6b063017 Add example -L kernel for algorithms with lookup table 2017-07-12 14:45:12 +02:00
jsteube
86e87f8957 Add example -L kernel for 64 bit based algorithms 2017-07-12 13:59:22 +02:00
jsteube
54b7505473 Add host modifications for -a 0 in combination with -L and modify an example kernel 2017-07-12 13:00:04 +02:00
jsteube
2c79d26778 Add -m 10700 pure kernel for -L support 2017-07-11 10:43:18 +02:00
jsteube
97390a9332 Fix -m 10700 if used on CPU 2017-07-10 12:19:04 +02:00
jsteube
093cf9af42 Fix datatype used in sha384_hmac_init_global_swap() 2017-07-10 12:15:38 +02:00
jsteube
28de23ec3e Simplify -m 10700 a bit 2017-07-10 12:10:49 +02:00
Jens Steube
e5a59a6611 Fix SIMD issue 2017-07-10 12:05:37 +02:00
jsteube
e70cc986da Small fix for SHA384 includes 2017-07-10 12:00:17 +02:00
jsteube
8a6e3a5275 Add support in HMAC for passwords larger than block size of the underlaying hash 2017-07-10 11:15:15 +02:00
jsteube
f619811b70 Remove PBKDF2-HMAC-MD5 includes password length limit 2017-07-09 23:53:53 +02:00
jsteube
97020f6521 Vectorized Ethereum Wallet + SCRYPT and added support for long passwords 2017-07-09 23:32:44 +02:00
jsteube
a91d048c04 Vectorized Ethereum Wallet, PBKDF2-HMAC-SHA256 and added support for long passwords 2017-07-09 23:10:28 +02:00
jsteube
25fba33901 Vectorized DPAPI masterkey file v1 and v2 and added support for long passwords 2017-07-09 23:05:14 +02:00
jsteube
32329cf3f4 Vectorized Juniper/NetBSD sha1crypt and added support for long passwords 2017-07-09 20:01:45 +02:00
jsteube
0da85fc1fd Vectorized iTunes backup >= 10.0 and added support for long passwords 2017-07-09 19:50:52 +02:00
jsteube
de9d026bb0 Vectorized iTunes backup < 10.0 and added support for long passwords 2017-07-09 19:37:36 +02:00
jsteube
8f73d356f2 Vectorized LUKS and added support for long passwords 2017-07-09 19:24:34 +02:00
jsteube
b1a88da83e Vectorized WinZip and added support for long passwords 2017-07-09 18:20:28 +02:00
jsteube
1049fa386a Add OPTI_TYPE_SLOW_HASH_SIMD_LOOP in interface.c where it was missing 2017-07-09 18:01:55 +02:00
jsteube
3141c14b0f Refactor OpenCL kernels to use normalized AES functions from inc_cipher_aes.cl 2017-07-09 17:39:38 +02:00
jsteube
709cfa2e91 Added long passwords support for KeePass 1 (AES/Twofish) and KeePass 2 (AES) 2017-07-09 15:12:11 +02:00
jsteube
837b5a31d1 Added long passwords support for AxCrypt 2017-07-09 00:13:00 +02:00
jsteube
dd14b798c3 Vectorized WPA/WPA2 PMK OpenCL kernel 2017-07-08 23:57:27 +02:00
jsteube
fc32b24236 Vectorized RAR5 and added support for long passwords 2017-07-08 23:55:56 +02:00
jsteube
da6bfc130e Fix invalid const keyword in OpenCL kernel function header 2017-07-08 23:45:00 +02:00
jsteube
933fa47d21 Vectorized Android FDE (Samsung DEK) and added support for long passwords 2017-07-08 21:56:36 +02:00
jsteube
edf904f309 Vectorized MS-AzureSync PBKDF2-HMAC-SHA256 and added support for long passwords 2017-07-08 21:39:15 +02:00
jsteube
af46a1560b Vectorized Blockchain, My Wallet and added support for long passwords 2017-07-07 23:32:41 +02:00
jsteube
a1321d2d64 Added long passwords support for BSDi Crypt, Extended DES 2017-07-07 22:16:42 +02:00
jsteube
02ce227ff1 Vectorized Oracle T: Type (Oracle 12+) and added support for long passwords 2017-07-07 22:09:51 +02:00
jsteube
6e57aa1c0f Vectorized eCryptfs and added support for long passwords 2017-07-07 21:46:41 +02:00
jsteube
729c5f09bc Vectorized PBKDF2-HMAC-SHA1 and added support for long passwords 2017-07-07 21:37:22 +02:00
jsteube
eda88e6c84 Vectorized PBKDF2-HMAC-MD5 and added support for long passwords 2017-07-07 16:58:28 +02:00
jsteube
d3e6ae42f0 Added long passwords support for 7-Zip 2017-07-07 16:48:18 +02:00
jsteube
27a57383f0 Vectorized Password Safe v3 and added support for long passwords 2017-07-07 16:02:49 +02:00
jsteube
0fae961111 Vectorized PBKDF2-HMAC-SHA256 and added support for long passwords 2017-07-07 15:44:05 +02:00
jsteube
e455561f77 Add -L support for -m 10800 in combination with -a 3 2017-07-07 12:48:35 +02:00
jsteube
5de48182b4 Fixed max password length limit in mode 10500 2017-07-07 12:33:06 +02:00
jsteube
bedc481390 Added long passwords support for SAP CODVN H (PWDSALTEDHASH) iSSHA-1 2017-07-07 12:14:06 +02:00
jsteube
f4301c9c22 Optimize use of AES128 and AES256 in Office 2007 2017-07-07 11:57:50 +02:00
jsteube
e2bc2a54c8 Optimize use of AES256 in Office 2013 2017-07-07 11:41:12 +02:00
jsteube
99dfdf466d Optimize use of AES128 in Office 2010 2017-07-07 11:35:47 +02:00
jsteube
ced2608326 Optimize Office 2007 and 2010 OpenCL _loop kernel 2017-07-07 10:42:39 +02:00
jsteube
8916de538a Vectorized MS Office 2013 and added support for long passwords 2017-07-07 10:38:05 +02:00
jsteube
51470b2b04 Vectorized MS Office 2010 and added support for long passwords 2017-07-07 10:03:59 +02:00
jsteube
61f39b37d2 Vectorized MS Office 2007 and added support for long passwords 2017-07-07 00:29:05 +02:00
jsteube
17b003b355 Vectorized Lotus Notes/Domino 8 and added support for long passwords 2017-07-06 14:57:28 +02:00
jsteube
df3890b49d Added long passwords support for SCRYPT 2017-07-06 14:27:36 +02:00
jsteube
ccd85f345d Vectorized 1Password, cloudkeychain and added support for long passwords 2017-07-06 14:11:33 +02:00
jsteube
d63e5f259f Remove some old code from -m 7900 2017-07-06 11:04:05 +02:00
jsteube
6cbd2acd24 Added long passwords support for Drupal7 2017-07-06 11:02:43 +02:00
jsteube
819b53eb1d Added long passwords support for sha256crypt $, SHA256 (Unix) 2017-07-05 13:43:14 +02:00
jsteube
7fec4f27d8 Vectorized OSX v10.8+ (PBKDF2-SHA512) and added support for long passwords 2017-07-05 13:01:55 +02:00
jsteube
b323682185 Merge branch 'master' of https://github.com/hashcat/hashcat 2017-07-05 12:54:21 +02:00
jsteube
f8cae33435 Fix some kernel headers 2017-07-05 12:54:15 +02:00
philsmd
03f4e2b3dc minor typo fixed in comment for the new update() functions 2017-07-05 12:16:37 +02:00
jsteube
195e3c744c Vectorized TrueCrypt PBKDF2-HMAC-Whirlpool and added support for long passwords 2017-07-05 10:08:47 +02:00
jsteube
81c9e3eb4f Backport 1024 and 1536 bit kernel for refactored -m 6212 2017-07-05 09:44:46 +02:00
jsteube
5105bc95a7 Add missing -m 6100 pure kernel for -a 3 2017-07-05 09:44:13 +02:00
jsteube
3c2e2c505e Add -L support for -m 6100 in combination with -a 3 2017-07-04 21:47:01 +02:00
jsteube
f2067d6962 Vectorized TrueCrypt PBKDF2-HMAC-RipeMD160 and added support for long passwords 2017-07-04 18:51:02 +02:00
jsteube
55874ec853 Vectorized VeraCrypt PBKDF2-HMAC-SHA256 and added support for long passwords 2017-07-04 15:40:34 +02:00
jsteube
907b065e00 Vectorized TrueCrypt PBKDF2-HMAC-SHA512 and added support for long passwords 2017-07-04 12:19:40 +02:00
jsteube
c9e98e48d3 Added long passwords support for Samsung Android Password/PIN 2017-07-04 11:49:43 +02:00
jsteube
5eb76ccdde Vectorized Password Safe v3 and added support for long passwords 2017-07-04 11:22:48 +02:00
jsteube
fc100a852b Added hash-mode 2501 = WPA/WPA2 PMK
Fixes https://github.com/hashcat/hashcat/issues/1287
Limited hash-mode 2500 to max length 63
Fixes https://github.com/hashcat/hashcat/issues/1286
2017-07-03 16:11:57 +02:00
jsteube
91f7acbde3 Remove more unused functions after refactor of -m 6x00 kernels 2017-07-03 13:20:05 +02:00
jsteube
1dfdefae69 Vectorized LastPass + LastPass sniffed kernel and added support for long passwords 2017-07-03 13:14:55 +02:00
jsteube
c3f0bb77dd Vectorized AIX {ssha1} kernel and added support for long passwords 2017-07-03 13:06:59 +02:00
jsteube
2e78cf1d58 Vectorized 1Password, agilekeychain kernel and added support for long passwords 2017-07-03 13:00:42 +02:00
jsteube
a8a1fe1b4f Vectorized AIX {ssha512} kernel and added support for long passwords 2017-07-03 12:44:03 +02:00
jsteube
bb1341015f Vectorized AIX {ssha256} kernel and added support for long passwords 2017-07-03 12:08:45 +02:00
jsteube
35e1ee6612 Functions append_0x02() no longer required after rewrite of -m 8800 2017-07-03 11:31:05 +02:00
jsteube
ccc9e46508 Vectorized Android FDE <= 4.3 kernel and added support for passwords up to length 256 2017-07-03 11:29:32 +02:00
jsteube
cbeb9c6e0c Remove unused functions from -m 2500 2017-07-03 11:22:50 +02:00
jsteube
d806aab2c3 Remove some unused functions 2017-07-03 11:18:40 +02:00
jsteube
b03382b334 Get rid of old truncate_block() 2017-07-03 11:07:37 +02:00
jsteube
64704f36c2 Add pure kernel for -m 6300 2017-07-03 10:50:27 +02:00
jsteube
34c85a659d Refactor -m 500 and -m 1600 to use new truncate_block() functions 2017-07-03 10:47:04 +02:00
jsteube
a22da36a00 Add different code generators for truncate_block(), add results to inc_common.cl and make use of them in m01800-pure.cl 2017-07-03 10:41:09 +02:00
jsteube
b9b2112b64 Add pure kernel for -m 1800 2017-07-02 23:27:54 +02:00
jsteube
8e1759650b Add -L support for -m 1700 in combination with -a 3 2017-07-02 19:05:38 +02:00
jsteube
58a66cf31d Fix function declaration of switch_buffer_by_offset_8x4_carry_be_S() 2017-07-02 18:39:55 +02:00
jsteube
a009f239d5 Backport generated scalar code to vector code 2017-07-02 18:26:44 +02:00
jsteube
fbfe81a0a1 Replace code in switch_buffer_by_offset_carry_be_S() with code generated with code generators 2017-07-02 18:15:13 +02:00
jsteube
b07b73f525 Add switch_buffer_by_offset_8x4_be_S() to inc_common.cl 2017-07-02 17:36:15 +02:00
jsteube
6feb0a1630 Rename switch_buffer_by_offset_64x1_le_S() to switch_buffer_by_offset_1x64_le_S() 2017-07-02 15:29:04 +02:00
jsteube
2a50c7ba61 Remove debugging line 2017-07-02 15:21:58 +02:00
jsteube
cd5223eb2f Replace code in switch_buffer_by_offset_be_S() with code generated with code generators 2017-07-02 15:18:44 +02:00
jsteube
521ece537c Replace code in switch_buffer_by_offset_64x1_le_S() in amp_a1.cl with generated code from code generators 2017-07-02 15:08:01 +02:00
jsteube
62d695d572 Rename switch_buffer_by_offset_le_S() to switch_buffer_by_offset_64x1_le_S() in amp_a1.cl 2017-07-02 15:05:55 +02:00
jsteube
48ce6cb71d Add append_0x80_8x4_S() and replace code in switch_buffer_by_offset_le_S() with generated code from code generators 2017-07-02 15:05:08 +02:00
jsteube
f5dca399ad Add -L support for -m 1600 2017-07-01 18:40:29 +02:00
jsteube
b0d5995689 Backport changes from inc_hash_md5.cl to inc_hash_md4.cl 2017-07-01 18:09:05 +02:00
Jens Steube
56dc8ae359 Add two functions md5_update_global_utf16le_swap() and md5_update_global_swap() for later use 2017-07-01 15:06:17 +02:00
jsteube
165380c454 Simplify WPA/WPA2 cracking kernel 2017-07-01 14:41:53 +02:00
jsteube
52c1e15f3f Move kernel-code for -L to standalone files with -pure suffix 2017-07-01 13:02:07 +02:00
jsteube
194af74e91 Add support for maximum bcrypt password length 2017-07-01 11:04:59 +02:00
jsteube
7914e075f6 This patch is an example of how to modify a fast -a 3 kernel to support password lengths up to 256 2017-06-30 17:21:30 +02:00
jsteube
cefd2ddb94 Tune AMD unroll settings for AMD-GPU-PRO 17.40 2017-06-29 15:50:46 +02:00
jsteube
f7a8e7c54b Multiple changes:
* Added more preparations to support to crack passwords and salts up to length 256
* Added option --length-limit-disable to disable optimization based on password- and salt-length
* Added option --self-test-disable to disable self-test functionality on startup
2017-06-29 12:19:05 +02:00
jsteube
c918173fcf Get rid of comb_t which can be safely replace with pw_t now 2017-06-25 00:56:25 +02:00
jsteube
045ac7d8e7 Modify amp_a1 to work with password length 256 2017-06-25 00:42:53 +02:00
Jens Steube
c2a770631f Merge pull request #1284 from neheb/master
Fix signed overflow warnings
2017-06-24 10:51:42 +02:00
Rosen Penev
2f3171fd98
Fix signed overflow warnings 2017-06-23 21:44:50 -07:00
jsteube
c59432a760 Add hcstat2 support to enable masks of length up to 256, also adds a filetype header 2017-06-23 12:13:51 +02:00
jsteube
120cf1d1ba Removed some unused functions, added -m 500 kernel with length 256 support but not activated because too slow 2017-06-23 09:24:50 +02:00
jsteube
71d4926afa Converted -m 400 to password length 256 support
Something weird happend here, read on!

I've expected some performance drop because this algorithm is using the password data itself inside the iteration loop.
That is different to PBKDF2, which I've converted in mode 2100 before and which did not show any performance as expected.

So after I've finished converting this kernel and testing everything works using the unit test, I did some benchmarks to see how much the
performance drop is.

On my 750ti, the speed dropped (minimal) from 981kH/s -> 948kH/s, that's mostly because of the SIMD support i had to drop.
If I'd turn off the SIMD support in the original, the drop would be even less, that us 967kH/s -> 948kH/s which is a bit of a more reasable
comparison in case we just want to rate the drop that is actually caused by the code change itself.

The drop was acceptable for me, so I've decided to check on my GTX1080.Now the weird thing: The performance increased from 6619kH/s to
7134kH/s!!

When I gave it a second thought, it turned out that:

1. The GTX1080 is a scalar GPU so it wont suffer from the drop of the SIMD code as the 750ti did
2. There's a change in how the global data (password) is read into the registers, it reads only that amount of data it actually needs by using
the pw_len information
3. I've added a barrier for CLK_GLOBAL_MEM_FENCE as it turned out to increase the performance in the 750ti

Note that this kernel is now branched into password length < 40 and larger.

There's a large drop on performance where SIMD is really important, for example CPU.

We could workaround this issue by sticking to SIMD inside the length < 40 branch, but I don't know yet how this can be done efficiently.
2017-06-22 13:49:15 +02:00
Jens Steube
0787b91327 Merge branch 'master' of https://github.com/hashcat/hashcat 2017-06-22 10:50:20 +02:00
Jens Steube
45b14ebf1c While not required now, it's better to use scalar functions in amplifier kernel in case it get's changed in the future 2017-06-22 10:28:35 +02:00
jsteube
cea78024bf Fix -m 2100 cracking if (password length & 31) == 0 2017-06-21 16:21:12 +02:00
jsteube
ad242c2f12 Working example of generic salt up to length 256 for mode 2100 2017-06-20 17:17:13 +02:00
jsteube
2c92465036 Add HMAC functionality to inc_hash_md4.cl and make DCC2 use it, resulting in support for longer domain names 2017-06-19 16:42:21 +02:00
jsteube
4174f06008 PoC using a length-independant MD4 hash processing in -m 2100 2017-06-18 23:31:40 +02:00
jsteube
a673aee037 Very hot commit, continue reading here:
This is a test commit using buffers large enough to handle both passwords and salts up to length 256.
It requires changes to the kernel code, which is not included in here.
It also requires some of the host code to be modified. Before we're going to modify kernel code to support the larger lengths I want to be
sure of:
1. Host code modification is ok (no overflows or underflows)
2. Passwords and Salts are printed correctly to status, outfile, show, left, etc.
3. Performance does not change (or only very minimal)
This is not a patch that supports actual cracking both passwords and salts up to length 256, but it can not fail anyway.
If if it does, there's no reason to continue to add support for both passwords and salts up to length 256.
2017-06-17 17:57:30 +02:00
jsteube
c9caca2b0c Increase max password length for DCC2 to 32 2017-06-16 19:28:55 +02:00
Jens Steube
bd01228ad5 Get rid of reqd_work_group_size(), no longer needed 2017-06-16 13:33:00 +02:00
jsteube
8a3ed7fe75 Small WPA improvement, do not check the same candidate twice for LE and BE 2017-06-16 11:33:49 +02:00
jsteube
cf57365e7c Check hashes_shown[] whenever calling mark_hash directly.
Not really needed right now (because those algorithms to have a single digest per salt), but this can change in the future
2017-06-16 10:48:10 +02:00
jsteube
4b2d9f0f29 Fix for https://github.com/hashcat/hashcat/issues/1276 2017-06-16 10:26:15 +02:00
jsteube
e9c010115d Fix some spacing 2017-06-16 10:25:36 +02:00
Jens Steube
e87fb31d3f WPA cracking: Improved nonce-error-corrections mode to use a both positive and negative corrections 2017-06-12 09:59:15 +02:00
Jens Steube
7e5b8d3f25 Added hash-mode 15500 = JKS Java Key Store Private Keys (SHA1) 2017-06-09 09:56:06 +02:00
jsteube
b8ad89c529 Rename function and variables containing 'unicode' to 'utf16le' because that's what's meant actually 2017-06-05 12:15:28 +02:00
Royce Williams
3fc185a66b tidy changes.txt and name normalizations 2017-06-04 13:54:41 -08:00
Jens Steube
542f73eb17 Move luks_tmp_t to correct position in inc_types.h 2017-06-04 21:52:35 +02:00
Jens Steube
9a1951d61c synchronize host and opencl types 2017-06-04 21:51:19 +02:00
Jens Steube
ae5fdba20f Add make_unicode and undo_unicode BE version 2017-06-04 21:16:41 +02:00
Jens Steube
b23ab71d5c - Added hash-mode 15600 = Ethereum Wallet, PBKDF2-HMAC-SHA256
- Added hash-mode 15700 = Ethereum Wallet, PBKDF2-SCRYPT
Resolves https://github.com/hashcat/hashcat/issues/1227
Resolves https://github.com/hashcat/hashcat/issues/1228
2017-06-03 23:23:03 +02:00
DoZ10
6ced398c3c Addressed comments and added 15400 to benchmark.c and tab_completion 2017-05-17 07:35:56 -04:00
DoZ10
5683df2e17 Fixed conflicts 2017-05-16 20:36:55 -04:00
DoZ10
8b6120243d Applied performance changes and fixed multi-mode bad implementation 2017-05-16 19:59:46 -04:00
DoZ10
f31f057113 Fixed minor error in offset calculation. 2017-05-15 22:49:44 -04:00
DoZ10
8dfd1bf066 Final. Implemented offset parameter to reach next keystream in kernels. Tested all kernels with scalar and vector modes 2017-05-15 18:34:34 -04:00
DoZ10
cfc3fa64c0 Implemented offset parameter to reach full ks block of 64 bytes 2017-05-15 08:47:40 -04:00
Fist0urs
a78dce94db All remarks treated:
1) done + got rid of all u8 datatypes in shared struct
	2) cf. previous
	3) necessary as this is computed in _init then used in _comp
	4) done
	5) done
	6) done => switch to 16
	7) done
2017-05-14 19:45:35 +02:00
DoZ10
5ab5e6c7b6 Added salt section comment in kernels 2017-05-14 07:17:51 -04:00
DoZ10
9dee1d274d Removed plain_length parameter and copied esalt buffer to salt ofr sorting mechanism. 2017-05-14 07:14:57 -04:00
DoZ10
0d3b5393ef Swapped mode 670 -> 15400 2017-05-14 06:52:14 -04:00
Fist0urs
7ff09c6710 Preparing PR 2017-05-09 20:14:07 +02:00
DoZ10
c50e8bc486 Fixed position parameters. Tested all kernels. Ok. 2017-05-07 14:02:00 -04:00
DoZ10
2fd31ed89f Completed kernel a3 2017-05-06 21:18:15 -04:00
DoZ10
dd1deb8ed3 Completed kernel a1 2017-05-06 21:12:53 -04:00
DoZ10
8511d9f047 Completed kernel a0. 2017-05-06 21:03:14 -04:00
DoZ10
3c67e0054c Implemented Perl test and fixed issues. Now have a working base. 2017-05-06 20:40:10 -04:00
DoZ10
cd9dc989ce Implemented Chacha20 crypto. 2017-05-06 14:24:50 -04:00
Fist0urs
40bbb0023c Merge branch 'master' of https://github.com/hashcat/hashcat into DPAPImk 2017-05-06 17:00:17 +02:00
Fist0urs
d537712f27 Both DPAPImk v1 and v2 work for single hash, still a bug on multi-hash
Remaining this bug + tests.pl before PR
2017-05-06 16:55:36 +02:00
DoZ10
f6cd42352d CPU parsing mostly done. Kernel showing good values. 2017-05-05 09:02:18 -04:00
DoZ10
152f0b5152 Init work on Chacha20 2017-05-04 22:34:52 -04:00
Jens Steube
b924901bb0 Merge pull request #1226 from DoZ10/master
Blake2b raw hash implementation
2017-05-03 16:38:03 +02:00
Fist0urs
29d331ee17 hmac-sha1 + SID almost working, padding problem 2017-05-02 23:56:00 +02:00
DoZ10
dcd8306b6f Fixed naming convention for kernel blake2_t type 2017-05-01 20:12:20 -04:00
DoZ10
ee558c625d Copied esalt buffer into kernel in the outer loop 2017-05-01 17:28:10 -04:00
DoZ10
8aa389b286 Fixed coding style convention for for() loops 2017-05-01 16:47:54 -04:00
Fist0urs
014278ab0e Working:
- MD4/sha1
  - hmac-sha1
  - pbkdf2-hmac-sha1
  - pbkdf2-hmac-sha512
Remaining:
  - handling of long salt (SID)
  - (AES256/DES3) + the end
2017-05-01 22:21:54 +02:00
DoZ10
22b9f80531 Inserted blake2b_sigma into kernel for perf gain. Standardized naming convention to blake2b_transform() 2017-05-01 16:19:20 -04:00
DoZ10
f0f96140b2 Fixed Attack modes 1 & 3 2017-04-30 21:10:54 -04:00
DoZ10
87e0281237 Moved init params in CPU (interface.c). Fixed vector-type problem in kernel a0. 2017-04-30 20:34:01 -04:00
Fist0urs
73d48dcd26 Initial commit, new format DPAPImk, works till hmac-sha1 2017-05-01 00:48:09 +02:00
mhasbini
5734741392 Add support for rule: eX 2017-04-30 16:23:39 +03:00
DoZ10
e71c68e0af Fixed endianness to remove unecessary kernel swaps ops 2017-04-29 22:10:06 -04:00
Jens Steube
0fcf51dee3 Fixed a missing type specifier in a function declaration of the RACF kernel 2017-04-27 17:56:38 +02:00
DoZ10
c1f8204b06 Prepared transform routine 2017-04-27 01:26:23 -04:00
DoZ10
0e018c717d Unrolled for() loops and removed S & P structs 2017-04-26 21:17:29 -04:00
DoZ10
76e3c0618e Fixed for() loops coding style 2017-04-26 16:15:50 -04:00
DoZ10
903e716b9a Swapped to outlen and tested. Okay. Still having problems with VECT_SIZE > 1 2017-04-25 18:28:40 -04:00
DoZ10
58c1f46b19 Merge and conflict resolve. 2017-04-24 21:47:35 -04:00
DoZ10
10629190e3 Now compiles for test.sh -m 600 -a 0 -V1 and -V2, but does not resolve on V2. 2017-04-23 20:32:34 -04:00
DoZ10
de477580a0 Code Cleanup and short doc additions 2017-04-22 21:50:22 -04:00
DoZ10
1386d0eecc Added kernel m00600_a1.cl 2017-04-22 10:44:38 -04:00
DoZ10
bb61408e89 Added kernel m00600_a0 2017-04-22 09:51:39 -04:00
DoZ10
60afdc2a30 Added function m04 in attack mode 3 2017-04-22 08:32:05 -04:00
DoZ10
72724ccba4 Removed rotr64_w() function. 2017-04-21 21:08:54 -04:00
DoZ10
b61d74255b Removed load64() and load64_reverse() functions. 2017-04-21 20:52:34 -04:00
DoZ10
ad305308bb Remove useless loop. 2017-04-20 19:32:54 -04:00
DoZ10
5689892e46 Removed useless spaces. 2017-04-20 19:30:04 -04:00
DoZ10
976f50a56b Code cleanup. 2017-04-20 19:24:28 -04:00
DoZ10
3ce9597685 Renamed load64_inv -> load64_reverse 2017-04-20 19:21:47 -04:00
DoZ10
afdef5ce1c fix: Kernel m00600_a3.cl function s04. 2017-04-19 13:27:28 -04:00
Jens Steube
a93a143d1e Use scalar swap32() version for scalar variable 2017-04-17 17:46:27 +02:00
DoZ10
d455c18d4b saved. 2017-04-16 11:51:59 -04:00
DoZ10
c3b95db072 rm: cleanup files. 2017-04-16 11:33:13 -04:00
DoZ10
d5ca5d59db add: lots of things.... 2017-04-16 10:07:12 -04:00
DoZ10
69dad31a29 fix: changed -m 6200 -> 600 2017-04-13 20:45:48 -04:00
DoZ10
abb0f84985 fix: m number in cl file. 2017-04-13 17:37:17 -04:00
DoZ10
bf45f87d39 random work... 2017-04-13 17:06:48 -04:00
DoZ10
4e9bb8b093 init 2017-04-13 08:38:59 -04:00
jsteube
26949a4fce WPA cracking: Improved nonce-error-corrections mode to fix corrupt nonce generated on Big-Endian devices
WPA cracking: Reduced --nonce-error-corrections default from 16 to 8 to compensate speed drop caused due to Big-Endian fixes
Fixes https://github.com/hashcat/hashcat/issues/1221
2017-04-10 10:11:32 +02:00
jsteube
d1b2fa0b31 Added hash-mode 15100 = Juniper/NetBSD sha1crypt 2017-03-23 16:44:32 +01:00
jsteube
2bc65c2c4d A bit a different _comp kernel iteration for WPA which can have lots of handshakes 2017-03-10 19:54:00 +01:00
jsteube
bea0e52cdb Remove unused variable 2017-03-09 10:17:34 +01:00
jsteube
c04dd5c8b1 WPA cracking: Added support for WPA/WPA2 handshake AP nonce automatic error correction
WPA cracking: Added parameter --nonce-error-corrections to configure range of error correction
2017-03-07 14:41:58 +01:00
jsteube
33a043ec63 Refactored internal use of esalt from sync with number of salts to sync with number of digests
This enables a true N esalts per salt feature which is required for WPA/WPA2 handling
In case we need that for a future algorithm, just make sure to have a unique value in digest in hash parser.
Fixes https://github.com/hashcat/hashcat/issues/1158
2017-03-07 09:44:58 +01:00
jsteube
89f8739dde Fixed overflow in bcrypt kernel in expand_key() function 2017-03-03 15:26:59 +01:00
philsmd
35c1f731b8
osx: some more volatile are required for luks/tc 2017-03-03 10:49:40 +01:00
philsmd
fbb1f92d2b
osx: -m 3200 = bcrypt needs a volatile variable 2017-03-02 22:24:12 +01:00
jsteube
72edd17481 Workaround -m 9100 to run on AMDGPU-Pro 2017-02-28 11:34:20 +01:00
jsteube
a5b8a91d58 Small -m 9100 cleanups 2017-02-28 11:13:57 +01:00
Jens Steube
e82ce9243d Added support for filtering hccapx message pairs using --hccapx-message-pair 2017-02-27 12:09:49 +01:00
Jens Steube
0fc949ef69 Fixed cracking of PeopleSoft Token if salt length + password length is >= 128 byte 2017-02-27 11:06:53 +01:00
jsteube
2ece9742e1 Compress multiple newlines to one 2017-02-26 15:42:56 +01:00
jsteube
d0fa9d059d Remove some unused macros 2017-02-26 15:34:45 +01:00
jsteube
22be61b20d Remove aligned __constant datatypes from OpenCL kernel function declarations 2017-02-26 15:16:36 +01:00
jsteube
6401c58568 Align all the __constant buffers to workaround OpenCL JIT compiler errors in NV drivers 378.x 2017-02-26 15:04:53 +01:00
jsteube
3fb433de60 Remove __constant from OpenCL kernel function declarations 2017-02-26 14:57:26 +01:00
jsteube
ecb851ecda Testwise workaround for -m 9100 on 378.x 2017-02-26 14:49:29 +01:00
jsteube
c094f3b511 Workaround added for NVidia OpenCL runtime: RACF kernel requires EBCDIC lookup to be done on shared memory 2017-02-25 17:36:29 +01:00
Jens Steube
2dd8018915 Fix -m 4520 for salt length exactly 15 2017-02-25 12:53:08 +01:00
jsteube
7aab78fb52 Fix broken SHA512 in LUKS due to union refactoring in tmps variable 2017-02-25 12:19:47 +01:00
jsteube
ba5fb80a38 Refactor kernel declarations for Skip32 to enable OSX using it 2017-02-23 15:18:41 +01:00
Jens Steube
4a3c90dd3c Fixes https://github.com/hashcat/hashcat/issues/1123 2017-02-22 20:53:14 +01:00
jsteube
9b08d4af0f Update hccapx format to version 4 2017-02-21 20:07:18 +01:00
Jens Steube
dd55c1eb66 WPA: Changed format for outfile and potfile from essid:mac1:mac2 to hash:mac_ap:mac_sta:essid
Fixes https://github.com/hashcat/hashcat/issues/1113
2017-02-19 14:45:27 +01:00
jsteube
bbb4c74e85 OpenCL Kernel: Remove "static" keyword from function declarations; Causes older Intel OpenCL runtimes to fail compiling 2017-02-17 10:11:05 +01:00
jsteube
63b7321be2 Fixed WPA/WPA2 cracking in case eapol frame is >= 248 byte 2017-02-13 18:31:15 +01:00
jsteube
366f5133ac Workaround added for AMDGPU-Pro OpenCL runtime: RAR3 kernel require a volatile variable to work correctly 2017-02-13 11:36:16 +01:00
philsmd
227c0be95e
resolves #1018: added -m 3910 = md5(md5($pass).md5($salt)) 2017-02-12 18:12:17 +01:00
jsteube
717da7f4f2 Added support for loading hccapx files 2017-02-06 17:03:42 +01:00
philsmd
c18c2950c9
mv switch_buffer_by_offset_8x4_le_S and switch_buffer_by_offset_8x4_le_VV () into inc_common.cl 2017-02-03 11:32:19 +01:00
philsmd
15a9644f2a
fixes #943: added -m 15000 = FileZilla Server >= 0.9.55 2017-02-03 00:07:31 +01:00
jsteube
c1d88f3234 Fix updated switch_buffer_* functions to run on AMD 2017-02-02 13:11:59 +01:00