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