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