jsteube
188a9568ce
Replace double MAYBE_VOLATILE
2018-08-09 19:44:54 +02:00
jsteube
dad05d9f69
Testing: Workaround some AMD OpenCL runtime segmentation faults
2018-08-09 13:03:22 +02:00
jsteube
103fdf04a1
Fixed a invalid scalar datatype return value in hc_bytealign() where it should be a vector datatype return value
2018-08-09 11:00:08 +02:00
Jens Steube
1c280e4a6e
Small performance boost for bcrypt on CPU
2018-08-02 14:20:04 +02:00
Michael Sprecher
3a321c8dce
Added hash-mode 16900 = Ansible Vault
2018-08-01 19:44:30 +02:00
jsteube
fbf434146d
Add set_mark_1x4() and set_mark_1x4_S()
2018-07-28 18:03:18 +02:00
Jens Steube
14c444fd47
Replace c_append_helper_mini[] table with on-the-fly calculation in order to workaround compiler bugs in AMD OpenCL runtime
2018-07-28 13:59:55 +02:00
jsteube
88ebca40b8
Added hash-mode 16800 = WPA-PMKID-PBKDF2
...
Added hash-mode 16801 = WPA-PMKID-PMK
Renamed lot's of existing WPA related variables to WPA-EAPOL in order to distinguish them with WPA-PMKID variables
Renamed WPA/WPA2 to WPA-EAPOL-PBKDF2
Renamed WPA/WPA2 PMK to WPA-EAPOL-PMK
2018-07-25 16:46:06 +02:00
philsmd
2e1845ec11
fixes #1624 : increase esalt/nonce buffer to 1024 for -m 11400 = SIP
2018-07-23 15:51:39 +02:00
Jens Steube
a43d3ad176
Rename some hashcat specific OpenCL functions to avoid conflicts with existing OpenCL functions from OpenCL runtime
2018-07-22 12:20:20 +02:00
Jens Steube
02a2495349
Switched array pointer types in function declarations in order to be compatible with OpenCL 2.0
2018-07-22 11:47:42 +02:00
Jens Steube
0ab7ab9cec
OpenCL kernels: Removed the use of 'volatile' in inline assembly instructions where it is not needed
2018-07-21 12:29:22 +02:00
jsteube
c1622d6593
Fixed detection of AMD_GCN version in case the rocm driver is used
2018-07-21 11:52:54 +02:00
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
