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