Gabriele Gristina
038bce131f
Add Double sha256 kernel module
2019-07-20 21:06:46 +02:00
Gabriele Gristina
ed77af9207
fix bug with -a3 -O
2019-07-13 16:25:31 +02:00
Gabriele Gristina
806b00168c
fix wrong OPTS_TYPE, add remaining 4710 kernels
2019-07-13 15:36:16 +02:00
Gabriele Gristina
bbcb23b038
Add hash mode 4710 - sha256(md5(pass))
2019-07-13 14:43:13 +02:00
Gabriele Gristina
c85cdd6cc8
fix kernel 20600 build error, issue #2094
2019-07-11 18:01:54 +02:00
Gabriele Gristina
660da2da3d
fix OpenCL compiler warning
2019-06-21 21:26:42 +02:00
Jens Steube
3234e9d6b5
Some more ROCm performance tuning for -m 77xx
2019-06-20 16:16:56 +02:00
Jens Steube
316095c151
Some more ROCm performance tuning
2019-06-20 10:04:31 +02:00
Jeremi M Gosney
871df0b81b
add hash mode 20600 (oracle transportation manager)
2019-06-18 11:41:41 -05:00
Jens Steube
6ec52bd342
ROCm JiT learned how to use V_ADD3_U32 efficiently
2019-06-18 12:41:59 +02:00
Jens Steube
5e0eb288c9
Use __launch_bounds__ in CUDA as replacement for reqd_work_group_size() in OpenCL
2019-06-16 18:01:26 +02:00
Jens Steube
dbbdb7e5ac
WipZip cracking: Added two byte early reject, resulting in higher cracking speed
2019-06-16 11:41:42 +02:00
philsmd
98759fba95
pkzip: some more missing DECLSPEC found
...
DECLSPEC should be specified on each and every OpenCL kernel function (in general)
2019-06-07 20:14:15 +02:00
philsmd
a661728256
pkzip: for u32 use MAX_DATA / 4
2019-06-07 19:42:28 +02:00
philsmd
01a511b9dd
minor: some code formatting changes for PKZIP
2019-06-07 17:24:13 +02:00
philsmd
316b2952b5
PKZIP: improve decompression and allow up to 320KB data length
2019-06-07 15:52:37 +02:00
Jens Steube
5920bd7f78
Speed up -m 19300 in general
2019-06-06 15:02:22 +02:00
Jens Steube
49c56f713a
Get rid of m08, m16, s08, s16 kernels in -m 9700 and -m 9720 since maximum password length for old office documents is 15
2019-06-04 17:01:35 +02:00
Jens Steube
e999ae8737
Speed up -m 11500 in general
2019-06-04 12:15:34 +02:00
Jens Steube
b66602f5f9
Fix -m 16100 in optimized -a 3 mode
2019-06-04 11:20:32 +02:00
Jens Steube
da10700840
Merge branch 'master' of https://github.com/hashcat/hashcat
2019-06-04 10:52:44 +02:00
Jens Steube
026436e2bc
Speed up -m 15000 in optimized -a 3 mode
2019-06-04 10:52:28 +02:00
Jens Steube
f689532e4c
Move P-box initializer values to constant memory in -m 3200
2019-06-03 14:40:51 +02:00
Jens Steube
1670ab06fa
Speed up -m 7700 and -m 7701
2019-06-03 10:28:34 +02:00
Sein Coray
3365040bc1
fixed two bugs where pkzip hashes wouldn't be cracked
2019-05-24 22:11:51 +02:00
Sein Coray
215440e43c
adding support for mixed multi-file pkzip hashes with mode 17225
2019-05-20 22:25:29 +02:00
Sein Coray
e08fc096cd
adding support to 17230 kernel to allow compression types 0 and 8
2019-05-20 19:54:16 +02:00
Sein Coray
4cf4891d1b
fixed length check for code1/2 as data length also contains iv length of 12
2019-05-20 19:43:01 +02:00
Jens Steube
07d8e5ef19
Merge pull request #2039 from s3inlc/master
...
fixed license text in pkzip master key kernels
2019-05-20 19:20:23 +02:00
Sein Coray
aed1910205
fixed license text in pkzip master key kernels
2019-05-20 19:17:47 +02:00
Jens Steube
a2dee17fc5
Merge pull request #2038 from s3inlc/pkzip-fix-4
...
Fix uncompressed hash attack when being longer than MAX_LOCAL
2019-05-20 19:15:34 +02:00
Jens Steube
5cd17df313
Merge pull request #2036 from s3inlc/pkzip-fix-3
...
Fixed inflate check on very short pkzip hashes
2019-05-20 11:58:13 +02:00
Sein Coray
edcdf004a5
Fixing issue when uncompressed pkzip hash is longer than MAX_LOCAL
2019-05-20 08:53:43 +02:00
Sein Coray
29ae5369c0
fixed copy-paste issue and missing m kernels
2019-05-20 07:21:04 +02:00
Sein Coray
67af2cf926
fixed inflate check on very short pkzip hashes
2019-05-17 16:03:26 +02:00
Sein Coray
c80bfde8f2
fix issue with pkzip hashes which have a larger offset value to be printed correctly
2019-05-17 15:28:49 +02:00
Sein Coray
15cbaa0f59
adding pkzip stream cipher kernels 20500 and 20510
2019-05-17 14:11:22 +02:00
Sein Coray
cd7b3ed672
fixed size of tmp to be checked on static huffman inflate
2019-05-15 16:49:52 +02:00
Sein Coray
0ea676907a
Merge remote-tracking branch 'upstream/master'
2019-05-15 14:33:32 +02:00
Sein Coray
845878f607
added missing pre-load
2019-05-14 15:43:01 +02:00
Sein Coray
06554f5071
applied speed improvements to modes 17220 and 17230
2019-05-14 15:39:31 +02:00
Sein Coray
d9e5a86765
updated m17230 to be cuda compatible
2019-05-14 14:56:37 +02:00
Sein Coray
8d1e737d60
updated m17220 to be cuda compatible
2019-05-14 14:47:48 +02:00
Sein Coray
c9c7261a05
updated m17210 to be cuda compatible
2019-05-14 14:38:55 +02:00
Sein Coray
b29019ae75
updated m17200 to be cuda compatible
2019-05-14 14:21:41 +02:00
Sein Coray
e300fe0d63
fixes to avoid crashing on gpu on multiple inner loops
2019-05-14 13:41:40 +02:00
Jens Steube
153a8704e0
Fix some register type in inline assembly in some ROCM section
2019-05-14 13:03:40 +02:00
Jens Steube
85f7d50fb8
Merge pull request #2021 from philsmd/master
...
added support for $electrum$3 hashes (-m 16600)
2019-05-14 11:57:46 +02:00
Sein Coray
e39a9284e1
Merge remote-tracking branch 'upstream/master'
2019-05-14 11:48:44 +02:00
Jens Steube
51ddf52369
Initialize CUDA vector datatypes to zero
2019-05-13 16:23:28 +02:00
philsmd
07a1bdb12c
added support for $electrum$3 hashes (-m 16600)
2019-05-13 14:34:15 +02:00
Jens Steube
c07f9c19c7
Reorder the TC/VC/DC header checks
2019-05-13 13:04:59 +02:00
Jens Steube
e2da5c8d57
Some unrolling for SHA2 based algorithms
2019-05-12 12:38:23 +02:00
Jens Steube
fa9d073f9a
Manually unroll sha2 hashes
2019-05-11 23:15:58 +02:00
Jens Steube
3ca3d1cc60
Fix kernel_rules variable name
2019-05-11 14:34:10 +02:00
Jens Steube
7832c54452
Fix constant memory use of bfs_buf
2019-05-11 09:32:16 +02:00
Jens Steube
53be3e74a3
Rename some variable to avoid collisions
2019-05-10 13:22:40 +02:00
Jens Steube
46f737c5af
Use real constant memory on CUDA
2019-05-10 13:22:26 +02:00
Sein Coray
371991e079
included speed improvements and feedback from atom applied to all 172xx kernel variants
2019-05-10 12:50:03 +02:00
Sein Coray
e4d8e4a7ad
bring fork up-to-date
2019-05-09 17:23:59 +02:00
Jens Steube
ce20a5ab6b
Fix uint4 rotate in scrypt based kernels for CUDA
2019-05-09 16:55:48 +02:00
Jens Steube
82927c13c8
Get rid of uchar4 in -m 9100
2019-05-09 13:09:27 +02:00
Jens Steube
ec4d4218c0
Add some missing operators for vector types
2019-05-09 12:59:36 +02:00
Jens Steube
6db4ab7e60
Fix scrypt based algorithms to work on CUDA
2019-05-09 11:11:52 +02:00
Jens Steube
027af75a39
Fix rotate function names
2019-05-08 20:42:46 +02:00
Jens Steube
6b7d064118
Replace (u32x) (...) with make_u32x (...)
2019-05-08 15:21:22 +02:00
Jens Steube
54dd2ea300
Use same settings for vector datatypes in inc_types.h as seen in cuda SDK vector_types.h
2019-05-07 16:07:28 +02:00
Jens Steube
7e5356126c
Fix more use of LOCAL_VK and LOCAL_AS
2019-05-07 12:22:37 +02:00
Jens Steube
03b2d3fb69
Fix use of LOCAL_VK and LOCAL_AS in -m 3200
2019-05-07 12:08:54 +02:00
Jens Steube
8ff8c5d536
Add LOCAL_VK to make use of __shared__
2019-05-07 09:01:32 +02:00
Jens Steube
bbed0cd67a
Fix test.sh and bitsliced algos
2019-05-06 15:06:02 +02:00
Jens Steube
d0bd33c9d1
Rename CONSTANT_AS to CONSTANT_VK
2019-05-06 14:34:16 +02:00
Jens Steube
ec9925f3b1
Warnings self-check and autotune with CUDA
2019-05-04 21:52:00 +02:00
Jens Steube
5ee033673c
Disable name mangling in NVRTC's PTX output and more
2019-05-03 15:50:07 +02:00
Jens Steube
58213c81d6
Add vector datatypes operators
2019-04-26 22:07:56 +02:00
Jens Steube
6a32e8ef18
Fix ulong datatype on Windows x64
2019-04-26 14:11:13 +02:00
Jens Steube
d9cb5cf8df
Fix recursion in inc_common.cl
2019-04-26 14:03:57 +02:00
Jens Steube
3b7304c9d8
Fix recursion in inc_platform.cl
2019-04-26 14:01:14 +02:00
Jens Steube
89119bf24a
Add missing inc_platform.h include
2019-04-26 13:59:43 +02:00
Jens Steube
00e1e32492
Replace barrier() with SYNC_THREADS()
2019-04-26 13:34:07 +02:00
Jens Steube
9faba41848
Use nvrtc to compile PTX (resulting PTX not yet used)
2019-04-26 13:28:44 +02:00
Jens Steube
4045e60021
Add nvrtc wrapper for later use
2019-04-26 10:03:16 +02:00
Jens Steube
4b986de5fb
Prepare native CUDA hybrid integration
2019-04-25 14:45:17 +02:00
Jens Steube
c02083281f
Fix undefined-internal warning message on ROCM
2019-04-24 14:17:34 +02:00
Jens Steube
c5c79feaaa
More cam_feistel() optimization
2019-04-23 21:56:40 +02:00
Jens Steube
f49d3f92e9
Reduce cam_feistel() xor count
2019-04-23 16:51:16 +02:00
Jens Steube
f10d27b2c9
Get rid of extract_byte() in inc_cipher_twofish.cl
2019-04-23 15:16:42 +02:00
Jens Steube
bf4b1a8e02
Remove duplicate 'static' declaration specifier in -m 19800 and -m 19900
2019-04-22 17:55:00 +02:00
jsteube
b9aaaf7809
Move 198xx DiskCryptor to 200xx to not collide with Kerberos 5, etype 17, Pre-Auth
2019-04-20 19:41:37 +02:00
Jens Steube
17ab30b29f
Merge pull request #1991 from brandoncasaba/master
...
Add hash modes 19800 (krb5pa etype 17) and 19900 (krb5pa etype 18)
2019-04-20 19:10:33 +02:00
jsteube
70fc36bf01
Reorganize inc_common.cl and make better use of HAS_* macros
2019-04-20 11:25:34 +02:00
jsteube
75b92c1ab1
Use both LE and BE modes for nonce error correction if none or both modes are set
2019-04-18 22:50:53 +02:00
jsteube
5b97fe7514
Workaround volatile for -m 2500 and -m 2501 for macosx
2019-04-18 16:52:14 +02:00
Jens Steube
38c1029f2e
Need volatile for IRIS GPU on Mac OSX for -m 2500 and -m 2501
2019-04-17 13:21:35 +02:00
jsteube
b2fecc5828
Fix typo in m11800_a0-pure.cl
2019-04-17 12:28:01 +02:00
Brandon Chalk
0a4ce19915
Merge branch 'master' of https://github.com/hashcat/hashcat
2019-04-16 10:03:05 -07:00
jsteube
74e3ede391
Test optimization for kernel include in pure mode
2019-04-15 18:11:15 +02:00
jsteube
a3b6e6f1b0
Fix hc_rotl64_S() for AMD devices
2019-04-15 12:11:37 +02:00
jsteube
8e89617015
Workaround setting password length bug in ROCM and POCL for -m 10100
2019-04-15 11:14:02 +02:00
jsteube
4cdff67011
Fix typo in OpenCL/m11800_a3-optimized.cl
2019-04-14 20:04:35 +02:00
jsteube
59ec6ac2e1
Fix some typos
2019-04-14 18:07:00 +02:00
jsteube
a671d501aa
Optimize some xxx_hmac_init functions
2019-04-14 17:03:37 +02:00
jsteube
eaefbec3fa
Rewrite streebogs input-data-length-bug workaround to make it easier to understand
2019-04-14 16:06:02 +02:00
jsteube
c34f75fe2d
Optimize some xxx_hmac_final functions
2019-04-14 15:59:03 +02:00
jsteube
9708275ac4
Fix missing rename constant variable name for sbob512_sl64
2019-04-14 15:50:57 +02:00
Jens Steube
fef62acade
Merge pull request #1995 from f0cker/library_fix
...
Fix for library compilation failure due to multiple def of sbob_xx64
2019-04-14 09:53:09 +02:00
jsteube
7c6970dbdd
Remove hard-coded static keyword from OpenCL kernels
2019-04-13 18:46:19 +02:00
jsteube
b7cdca09c4
OpenCL Runtime: Workaround JiT compiler error on ROCM 2.3 driver if the 'inline' keyword is used in function declaration
2019-04-13 13:46:55 +02:00
jsteube
85d58b03e1
Fix some signed/unsigned integer comparison warnings
2019-04-10 12:23:39 +02:00
Brandon Chalk
61ac3e3282
Minor speed and readability improvements to mode 19800 and 19900
2019-04-09 23:25:01 -07:00
Brandon Chalk
95c74c52c1
Add hash modes 19800 (krb5pa etype 17) and 19900 (krb5pa etype 18)
2019-04-09 21:58:42 -07:00
jsteube
8f62085ea4
Add missing cipher cascade support for DiskCryptor
2019-04-08 19:33:24 +02:00
philsmd
a468249aba
diskcryptor: allow cd/iso encrypted images
2019-04-08 18:11:14 +02:00
jsteube
7bd3d55a3f
Add missing OpenCL kernel for -m 19812 and -m 19813
2019-04-08 11:07:12 +02:00
jsteube
4cd98603f5
Add DiskCryptor modules for 1024 bit and 1536 and support for Serpent and Twofish. Be careful, untested yet!
2019-04-08 09:31:31 +02:00
philsmd
60c255a69d
added -m 19800 = DiskCryptor AES
2019-04-07 18:08:56 +02:00
jsteube
c911f2bd83
Fix some OpenCL JiT compiler warnings on ROCM
2019-04-06 17:24:57 +02:00
f0cker
d2f29a2f21
Fix for library compilation failure due to multiple def of sbob_xx64, renamed these
2019-04-06 10:42:18 -04:00
jsteube
7e55aad703
Merge branch 'master' of https://github.com/hashcat/hashcat
2019-04-06 14:03:13 +02:00
jsteube
ed8af919a1
Use local memory for DES constants in -m 15300 and for AES in -m 137xx
2019-04-06 13:19:54 +02:00
jsteube
d706d19b4f
Fix some uninitialized variables
2019-04-05 22:25:28 +02:00
jsteube
ace765bf96
Fix 'E' rule in optimized mode if password is longer than 16 characters
2019-04-05 12:24:53 +02:00
jsteube
3274220d7f
Fix l32_from_64_S(), h32_from_64_S(), hl32_to_64_S() and h32_from_64() in inc_common.cl
2019-04-04 22:05:45 +02:00
jsteube
875718fb8e
Backport -m 500 pure kernel code to -m 6300 enables using it on AMDGPU driver
2019-04-04 21:05:07 +02:00
jsteube
cd88410781
Backport -m 500 pure kernel code to -m 1600 enables using it on AMDGPU driver
2019-04-04 21:02:26 +02:00
jsteube
d7d716f3ab
Make it easier to include OpenCL kernels into modules
2019-04-04 20:01:37 +02:00
jsteube
9ced13cc94
Get rid of CONSTSPEC macro in OpenCL kernels
2019-04-04 10:15:34 +02:00
jsteube
13edc32fb4
Fix some unused variable warnings
2019-04-03 21:53:34 +02:00
jsteube
b8d609ba16
WPA/WPA2 cracking: In the potfile, replace password with PMK in order to detect already cracked networks across all WPA modes
2019-04-02 11:24:22 +02:00
Jens Steube
4115e6b825
Update some unstable_warning on Intel CPU
2019-04-01 11:22:51 +02:00
jsteube
319bf80178
Fix hash_encode() salt position
2019-03-31 20:17:17 +02:00
jsteube
76bf5173ba
Do not use MAYBE_UNUSED in inc_common.cl
2019-03-31 19:45:15 +02:00
jsteube
1c45de4b63
Remove more old GCC diagnostic options
2019-03-31 19:41:28 +02:00
jsteube
d80603648e
Replace plain_t with void in module_build_plain_postprocess() enables some options
2019-03-30 16:32:11 +01:00
jsteube
9574862a19
Inline hc_byte_perm in hc_bytealign
2019-03-28 13:37:28 +01:00
jsteube
ffd8ec9001
Fixed output password of 'e' rule in pure and cpu rule engine if separator character is also the first letter
2019-03-28 13:07:39 +01:00
jsteube
5a1d929628
Fix some missing code changes after hc_bytealign() was changed
2019-03-28 12:26:24 +01:00
jsteube
ea7d9c50b6
Fix missing const keyword in some includes
2019-03-27 11:38:49 +01:00
jsteube
2f972c0d36
Use apply_rules_vect_optimized() instead of apply_rules_vect() in -m 15400
2019-03-27 11:17:30 +01:00
jsteube
19a77c59f9
Removed some unused includes from ChaCha20 combinator attack kernel
2019-03-27 11:17:13 +01:00
jsteube
4fe1a4b258
Use apply_rules_vect_optimized() instead of apply_rules_vect() in -m 600
2019-03-27 11:12:22 +01:00
jsteube
1a44581524
Removed some unused includes from Blake2 combinator attack kernel
2019-03-27 11:10:26 +01:00
Jens Steube
0fb3b3c83e
Declare internal functions in OpenCL kernels as static
2019-03-26 11:03:25 +01:00
jsteube
0a8c7fab1c
Do not use ulong. It causes the 32 bit compilation to fail
2019-03-25 21:04:21 +01:00
jsteube
ecbd4a51c8
Add remaining emu_inc_* sources to Makefile
2019-03-25 16:15:58 +01:00
jsteube
5b667d2c01
Get rid of src/cpu_aes.c, src/cpu_des.c, src/cpu_md5.c and src/cpu_sha256.c
2019-03-25 15:54:58 +01:00
jsteube
fb8a9d7c40
Get rid of cpu_sha1.c
2019-03-25 14:08:59 +01:00
Sein Coray
21a214e26a
Merge branch 'master' of https://github.com/hashcat/hashcat
2019-03-25 12:27:33 +01:00
jsteube
66d94b06e4
Get rid of src/rp_kernel_on_cpu.c and src/rp_kernel_on_cpu_optimized.c and use OpenCL emulated kernel version
2019-03-25 12:24:04 +01:00
Sein Coray
798f026e3a
adjusted kernels to newest refactoring
2019-03-25 10:57:18 +01:00
Sein Coray
959852a1aa
Merge branch 'master' of https://github.com/hashcat/hashcat
2019-03-25 10:47:00 +01:00
jsteube
13097fefc7
reorder functions for better overview
2019-03-24 15:23:11 +01:00
jsteube
194fd7e6d1
Fix some invalid code sections caused from conversion
2019-03-24 13:46:06 +01:00
jsteube
977199698f
Move some macros from .h to .cl sources
2019-03-24 00:26:10 +01:00
jsteube
e80b1838e8
Rename some functions in inc_common.cl to avoid conflicts with bitops.c
2019-03-23 22:15:38 +01:00
jsteube
adeeaee84a
Replace __kernel, __constant, __global and __local qualifiers with macro for better control
2019-03-22 22:27:58 +01:00
jsteube
340d2ae7d3
Add headers for OpenCL include files
2019-03-22 22:00:07 +01:00
Sein Coray
67994fc177
removed unecessary parts causing build on some Intel OpenCLs to fail
2019-03-22 21:09:38 +01:00
jsteube
0413314ae4
Fix missing include in amp_a0
2019-03-22 15:44:21 +01:00
jsteube
c9d60c079f
Prepare OpenCL kernels for non-static compilation
2019-03-22 15:16:25 +01:00
jsteube
5f5468be6f
Add missing DECLSPEC in OpenCL rule functions
2019-03-22 13:14:25 +01:00
jsteube
c1d5d2ff45
Improve grouping of include calls as a preparation
2019-03-22 12:58:56 +01:00
jsteube
305a044ec6
Remove some old function headers
2019-03-22 09:52:54 +01:00
jsteube
7d4bea41a0
Get rid of OpenCL/inc_hash_constants.h and OpenCL/inc_hash_functions.cl
2019-03-21 23:00:38 +01:00
Sein Coray
c768073f60
fixed static limit for uncompress output to length defined
2019-03-21 21:37:24 +01:00
Sein Coray
97249ba1ca
added pkzip kernels 17200, 17210, 17220 and 17230
2019-03-21 20:18:39 +01:00
jsteube
59ecdbd319
Add volatile to inline assembly instructions, it has no influence on cracking performance but compile time reduces
2019-03-20 11:05:34 +01:00
Jens Steube
af129d32da
Merge pull request #1941 from s3inlc/master
...
Added hash mode 19500
2019-03-19 10:44:06 +01:00
Sein Coray
836c91c8c2
Merge branch 'master' of https://github.com/s3inlc/hashcat
2019-03-18 17:47:42 +01:00
Sein Coray
17083b8f56
applied requested final changes
2019-03-18 17:47:35 +01:00
jsteube
e0f875463f
Fix HAS_VBFE use on AMDGPU driver
2019-03-18 12:48:59 +01:00
jsteube
5ef67a8ab7
Apply previous blowfish optimization for -m 3200 also on -m 9000 and -m 18600
2019-03-18 12:42:47 +01:00
Sein Coray
1143cef606
Merge branch 'master' into master
2019-03-18 09:48:09 +01:00
Sein Coray
7321b03102
Updated mode 19500 based on comments
2019-03-18 09:45:04 +01:00
jsteube
5ecbcde945
Cracking bcrypt: Use a feedback from the OpenCL runtime to dynamically find out optimal thread count
2019-03-17 14:17:35 +01:00
jsteube
9a70655f34
No more need for rcon[] in AES
2019-03-16 21:15:05 +01:00
jsteube
7cb510f1ce
More manually unrolled cipher code
2019-03-16 21:11:02 +01:00
jsteube
a063e9ef62
Remove automatic unrolling on some ciphers and replace with manually unrolled code
2019-03-16 19:00:36 +01:00
jsteube
70d1343d57
Fix variables s_Ch and s_Cl in whirlpool hashes in non REAL_SHM mode
2019-03-16 16:51:54 +01:00
jsteube
970e5f3518
Fix -m 6100 in optimized mode for use with REAL_SHM
2019-03-15 23:27:44 +01:00
jsteube
218322f630
Support for inline VeraCrypt PIM Brute-Force
2019-03-15 21:48:49 +01:00
Fist0urs
afb010870a
Add hash modes 19600 (krb5tgs enctype 17) and 19700 (krb5tgs enctype 18)
2019-03-13 17:20:04 +01:00
jsteube
26d7602bbd
No longer need reqd_work_group_size() in OpenCL kernels since modules taken care of this using threads_min and threads_max
2019-03-10 12:13:14 +01:00
jsteube
85452e12be
Fix inc_cipher_des.cl for vector use
2019-03-10 11:07:52 +01:00
jsteube
84d6b8ecc1
Add function prototypes in OpenCL kernels to make some compilers happy
2019-03-09 09:05:44 +01:00
jsteube
ce32d57f9b
Make use of inc_cipher_des.cl in -m 3100
2019-03-09 08:43:07 +01:00
jsteube
27610008cb
Testwise disable aligned constants
2019-03-09 08:11:29 +01:00
jsteube
aeb570ca8b
Add inc_cipher_des.cl and make use of SHM_TYPE
2019-03-09 08:07:50 +01:00
jsteube
8c47a947a1
Fix some includes in -m 14600
2019-03-08 15:39:32 +01:00
Jens Steube
bc8ffadccc
Remove more occurances of netntlm_t in Whirlpool kernel
2019-03-08 15:26:50 +01:00
Jens Steube
07b8cafcad
Fix datatypes in -m 14600
2019-03-08 13:53:34 +01:00
Jens Steube
d64529ed47
Fix kernel header in -m 14900
2019-03-08 13:41:20 +01:00
Jens Steube
fda04609cd
Fix datatypes in -m 14600
2019-03-08 13:37:13 +01:00
Jens Steube
898000b7aa
Fix datatypes in -m 12400
2019-03-08 13:17:07 +01:00
Jens Steube
09b8bfc3a5
Fix datatypes in -m 10700
2019-03-08 13:13:01 +01:00
Jens Steube
7932d8c91f
Fix datatypes in -m 64xx and -m 137xx
2019-03-08 13:10:58 +01:00
Jens Steube
34850df1c0
Fix datatypes in -m 2501 and -m 16801
2019-03-08 13:05:23 +01:00
Jens Steube
74abff44d9
Move remaining module specific typedefs and structures from inc_types.cl to kernels
2019-03-08 12:50:31 +01:00
Jens Steube
30681e5151
Move remaining module specific typedefs and structures from inc_types.cl to kernels
2019-03-08 11:14:33 +01:00
Jens Steube
4bce25dd9d
Move remaining module specific typedefs and structures from inc_types.cl to kernels
2019-03-08 10:18:20 +01:00
Jens Steube
cace799187
Move more no unrolling exceptions from inc_vendor.cl to specific modules
2019-03-07 18:57:04 +00:00
Sein Coray
ee97d18af9
Added hash mode 19500
2019-03-06 21:15:09 +01:00
Sein Coray
deacf3a2b9
fix constant name for sha224
2019-03-06 00:13:10 +01:00
jsteube
eb07d2108c
Fixed cracking of NetNTLMv1 passwords in mask-attack mode if mask > length 16 (optimized kernels only)
2019-03-05 05:14:34 +01:00
jsteube
8d6a69b2a9
Fix salt array for -m 19300 in -a 3 mode if SIMD is used
2019-03-04 13:31:49 +01:00
jsteube
a9bafb7edb
Add more valid plaintext pattern for -m 12700/15200
2019-03-04 09:28:57 +01:00
jsteube
dc9279c95c
New Strategy: Instead of using volatile, mark the mode as unstable. Remove all volatiles
2019-03-03 19:18:56 +01:00
jsteube
7c1120b784
Migrate unroll handling for -m 15900 from OpenCL/inc_vendor.cl to module_jit_build_options()
2019-03-03 13:44:25 +01:00
jsteube
ed7765d8fd
Manually unroll aes128_InvertKey and aes256_InvertKey to workaround some JiT compiler errors
2019-03-03 13:39:52 +01:00
jsteube
0fce6b663b
Testwise remove all hardcoded volatiles from cipher includes
2019-03-03 12:12:01 +01:00
jsteube
1714382daa
Migrate unroll handling for -m 14100 from OpenCL/inc_vendor.cl to module_jit_build_options()
2019-03-03 11:55:42 +01:00
jsteube
b0a83b28d1
Migrate unroll handling for -m 12300 from OpenCL/inc_vendor.cl to module_jit_build_options()
2019-03-03 11:47:33 +01:00
jsteube
1937b67595
Migrate unroll handling for -m 8200 from OpenCL/inc_vendor.cl to module_jit_build_options()
2019-03-03 11:43:47 +01:00
jsteube
fee37ac4c3
Migrate unroll handling for -m 8000 from OpenCL/inc_vendor.cl to module_jit_build_options()
2019-03-03 11:39:18 +01:00
jsteube
f92ebc6586
Fix -m 15200 and -m 12700 search string in kernel
2019-03-03 08:47:58 +01:00
jsteube
0b1169e523
My Wallet: Added additional plaintext pattern used in newer versions
2019-03-02 21:58:50 +01:00
jsteube
93760dab34
Fix -m 8100 in pure kernel mode for password length 256
2019-03-02 17:33:22 +01:00
jsteube
88a051629c
Support module specific JiT compiler build options
2019-03-02 11:12:13 +01:00
jsteube
a5743c5572
Fix invalid kernel declaration in optimized Whirlpool kernel
2019-03-01 21:45:25 +01:00
jsteube
4359418918
OpenCL Runtime: Not using amd_bytealign (amd_bitalign is fine) on AMDGPU driver drastically reduces JiT segfaults
2019-03-01 15:34:49 +01:00
jsteube
e1fe3e755b
Optimize some switch_buffer_* functions for generic OpenCL devices (CPU, various OSX, ...)
2019-03-01 14:49:00 +01:00
jsteube
23917455ef
Added hash-mode 19300 sha1(..)
2019-02-28 20:00:52 +01:00
jsteube
87c24200da
Added QNX /etc/shadow hash cracking support
...
- Added hash-mode 19000 QNX /etc/shadow (MD5)
- Added hash-mode 19100 QNX /etc/shadow (SHA256)
- Added hash-mode 19200 QNX /etc/shadow (SHA512)
Implements #35
2019-02-27 17:53:00 +01:00
jsteube
5da1e4b872
Fixed maximum password length limit which was announced as 256 but actually was 255
2019-02-26 21:20:07 +01:00
jsteube
b1a056f1b0
Add hc_lop_0x96 as template for eventual later use
2019-02-24 17:20:09 +01:00
jsteube
540b405e3a
Replace IS_ROCM with HAS_VPERM and HAS_VADD3
2019-02-24 10:12:48 +01:00
jsteube
2c0c82c8af
OpenCL Runtime: Workaround JiT compiler error on AMDGPU driver compiling WPA-EAPOL-PBKDF2 OpenCL kernel
2019-02-24 09:10:37 +01:00
jsteube
69c3ea2d27
Fixed invalid transfer from __constant to __local memory in -m 9100
2019-02-24 08:26:10 +01:00
jsteube
be365acef8
Remove some optimization in -m 18700 which ROCM doesn't like
2019-02-24 08:02:13 +01:00
jsteube
d0d4ce9f8c
Added hash-mode 18800 Blockchain, My Wallet, Second Password (SHA256)
2019-02-23 17:45:02 +01:00
jsteube
684256022e
Small fix for -m 18900
2019-02-22 16:00:08 +01:00
jsteube
29fedf2c41
Added hash-mode 18900 Android Backup
2019-02-22 15:49:47 +01:00
jsteube
3dd0a7140d
Respect combs_mode in -a 1/6/7 attack in -m 18700
2019-02-22 12:43:22 +01:00
jsteube
a0fba5fb11
Improve -m 18700 cracking speed
2019-02-22 12:33:16 +01:00
jsteube
b4d52e412b
Rename -m 18700 to Java Object hashCode()
2019-02-22 09:30:56 +01:00
jsteube
51eb9ebff7
Added hash-mode 18700 DJB 32
2019-02-21 13:52:01 +01:00
jsteube
9fc193ce47
Bitcoin Wallet: Be more user friendly by allowing a larger data range for ckey and public_key
2019-02-20 16:20:28 +01:00
jsteube
63fac132e3
Fix cracking streebog 256/512 hmac cracking with password length > 64
2019-02-19 17:17:01 +01:00
jsteube
f4e43da456
Fix whirlpool final() handling
2019-02-17 08:57:51 +01:00
jsteube
158b93832c
Fixed cracking of Cisco-PIX and Cisco-ASA MD5 passwords in mask-attack mode if mask > length 16
2019-02-15 15:50:58 +01:00
jsteube
dd293f7a93
Fixed -m 600 in -a 3 mode for passwords > 16
2019-02-13 13:48:31 +01:00
jsteube
e571b890e9
Fixed length check for raw PBKDF2 modules
2019-02-13 10:03:07 +01:00
jsteube
c16a3feabc
Add -m 16500 module
2019-02-11 13:11:51 +01:00
jsteube
57da64533c
Add missing kernel in -m 1100 -a 3 -O mode
2019-02-10 10:26:35 +01:00
jsteube
3d203af066
Add module for -m 5700 and -m 5800
2019-02-09 10:03:58 +01:00
jsteube
0bd244c051
Undef some macros to avoid collisions
2019-02-07 16:59:10 +01:00
jsteube
c88a837196
Rename d_scryptVX_buf to d_extraX_buf
2019-01-04 11:21:42 +01:00
R. Yushaev
393916c0bf
Allow cracking non-unique salts for Office 2013
...
With hash-mode 9600 (MS Office 2013) there can be multiple hashes with
the same salt but with different encryption verifiers in esalt_bufs.
This commit adds the functionality to execute _comp kernels for
different hashes after deriving their common key once.
Fixes #1826
2018-12-18 14:32:56 +01:00
R. Yushaev
b5a7e967c1
Add support for Open Document Format 1.1
...
Contains a kernel for the ODF 1.1 encryption implemented in OpenOffice.
The algorithm uses a SHA-1 checksum, a PBKDF2-HMAC-SHA1 key derivation
with 1024 iterations and Blowfish-CFB encryption.
Valid hashes can be extracted with the libreoffice2john.py script,
available from the John the Ripper Jumbo repository at
https://github.com/magnumripper/JohnTheRipper/blob/bleeding-jumbo/run/libreoffice2john.py
You have to remove the filename suffix at the end of the hash before
passing it to hashcat. Also see 'hashcat -m18600 --example-hashes'.
You can leave the filename prefix if you use the --username option to
process those hashes.
- Add hash-mode 18600 (Open Document Format (ODF) 1.1 (SHA-1, Blowfish))
- Tests: add hash-mode 18600 (Open Document Format (ODF) 1.1 (SHA-1, Blowfish))
2018-12-14 13:23:52 +01:00
Sein Coray
a70a0513bf
Added hash mode 18500 sha1(md5(md5($pass)))
...
closes hashcat/hashcat#1652
2018-12-10 16:11:11 +01:00
Jens Steube
15ece0902f
Merge pull request #1804 from Naufragous/odf-cracking
...
Add support for Open Document Format 1.2
2018-12-07 09:36:31 +01:00
Jens Steube
9d213147e8
Merge pull request #1805 from mcovalt/electrum_salt_type_2
...
Electrum Salt-Type 2
2018-12-07 09:33:13 +01:00
R. Yushaev
6a5b0c821e
Add support for Open Document Format 1.2
...
Contains a kernel for the latest ODF 1.2 encryption implemented in
LibreOffice. The algorithm uses a SHA-256 checksum, a PBKDF2-HMAC-SHA1
key derivation with 100000 iterations and key stretching and AES-CBC
encryption.
Valid hashes can be extracted with the libreoffice2john.py script,
available from the John the Ripper Jumbo repository at
https://github.com/magnumripper/JohnTheRipper/blob/bleeding-jumbo/run/libreoffice2john.py
You have to remove the filename suffix at the end of the hash before
passing it to hashcat. Also see 'hashcat -m18400 --example-hashes'.
You can leave the filename prefix if you use the --username option to
process those hashes.
- Add hash-mode 18400 (Open Document Format (ODF) 1.2 (SHA-256, AES))
- Tests: add hash-mode 18400 (Open Document Format (ODF) 1.2 (SHA-256, AES))
2018-12-06 18:00:09 +01:00
Matt Covalt
bb9328f48b
Add salt type 2 support for optimized functions
2018-12-05 10:42:23 -08:00
Matt Covalt
a913db6390
Add salt type 2 support for pure functions
2018-12-05 10:37:40 -08:00
Matt Covalt
36bdcf844a
Add function to check if vector represents a Base58 string
2018-12-05 10:20:41 -08:00
Sein Coray
c941e55a35
Extended IKE PSK md5/sha1 (-m 5300/5400) to print hashes correctly
2018-12-05 12:57:54 +01:00
R. Yushaev
5efebb7b48
Cleanup VeraCrypt related code
...
Remove unnecessary constant variables by hardcoring values instead of
looking up. Precalculate swaps that are known at compile time. Hardcode
hashes_shown offset as zero for all TC / VC kernels.
2018-11-29 12:49:03 +01:00
R. Yushaev
baf47d409e
Add Camellia support for VeraCrypt kernels
...
Adds suport for the Japanese cipher Camellia with 256-bit keys as used
by VeraCrypt.
- Add Camellia header decryption checks to all VeraCrypt kernels
- Add test containers for remaining cipher combinations
2018-11-28 14:21:14 +01:00
Jens Steube
65abccc93e
Migrate inc_truecrypt_keyboard.cl into inc_common.cl
2018-11-25 18:34:57 +01:00
Jens Steube
fca4f7e8a6
Prepare to use --keyboard-layout-mapping for algorithms other than TC/VC
2018-11-25 18:21:07 +01:00
Jens Steube
ee2854ec2a
Support multi-byte characters for TC/VC keyboard layout mapping tables
2018-11-25 13:31:37 +01:00
R. Yushaev
8b04be0e93
Add Kuznyechik support for VeraCrypt kernels
...
Adds support for the Russian cipher specified in GOST R 34.12-2015, also
known as Kuznyechik (Grasshopper).
- Add Kuznyechik header decryption checks to all VeraCrypt kernels
- Add test containers for available Kuznyechik cipher combinations
2018-11-22 16:07:45 +01:00
Jens Steube
c15f741dca
Fixed out-of-boundary read in DPAPI masterkey file v2 OpenCL kernel
2018-11-21 14:55:22 +01:00
Jens Steube
e117e750fc
Add restrict keyword to kernel declarations to help the compiler with caching optimizations
2018-11-21 13:00:30 +01:00
Jens Steube
64dfd40113
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 15:44:24 +01:00
Jens Steube
53c8600089
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 15:34:43 +01:00
Jens Steube
0e428b3c40
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 15:32:41 +01:00
Jens Steube
c9da60c73a
Fixed thread count maximum for pure kernels in straight attack mode
2018-11-20 15:29:24 +01:00
Jens Steube
2a6444c05a
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 15:26:46 +01:00
Jens Steube
ae577410d0
OpenCL Device: Do a real query on OpenCL local memory type instead of just assuming it
2018-11-20 10:06:34 +01:00
jsteube
240f6298be
Fix some leftovers from switching kernel parameters to macros
2018-11-17 23:14:12 +01:00
jsteube
eec1fba4c3
Fix some leftovers from switching kernel parameters to macros
2018-11-17 17:57:15 +01:00
jsteube
a930c5c24e
Do not use KERN_ATTR_VECTOR() if there is no SIMD
2018-11-16 23:35:58 +01:00
jsteube
c672182b44
Fix some leftovers from switching kernel parameters to macros
2018-11-16 23:04:20 +01:00
R. Yushaev
fbbe5f6282
Use macros in remaining kernel functions
...
The 7zip, scrypt and stdout kernels differ from the others in their
function declarations somewhat. Unify them and substitute with macros.
Also remove a few superfluous (bogus) consts which were introduced in
the previous PR.
2018-11-16 14:30:45 +01:00
Jens Steube
6d39fb1feb
Make all kernel parameter macro helper a function
2018-11-16 14:17:01 +01:00
R. Yushaev
5de004103a
Replace kernel parameter lists with macros
...
Substitute long parameter lists in ~2900 kernel function declarations
with macros. This cleans up the code, reduces probability of copy-paste
errors and highlights the differences between kernel functions. Also
reduces the size of the OpenCL folder by ~3 MB.
2018-11-16 11:44:33 +01:00
R. Yushaev
31dc7a3453
Add macros for kernel function declarations
...
Most of the kernel functions use nearly identical parameter lists.
Essentially, there are four parameters that vary, except for a dozen odd
kernel functions (e.g. stdout, zip, scrypt). This means that the function
declarations can be hidden behind a few simple macros with up to two
parameters.
2018-11-16 11:44:26 +01:00
R. Yushaev
b80ada1d65
Unify esalt_bufs parameter declarations
...
In preparation for the abstraction of long repetitive kernel function
declarations, rename the salt buffer pointers to *esalt_bufs. Also
declare them const where they are not.
2018-11-16 10:28:54 +01:00
R. Yushaev
3f0a3ef3f7
Unify kernel function parameter names
...
In preparation for the abstraction of long repetitive kernel function
declarations, adjust parameter names in a few deviating kernels.
2018-11-16 10:28:54 +01:00
Jens Steube
49fc7d45b7
- Keymaps: Added hashcat keyboard mapping us.hckmap (can be used as template)
...
- Keymaps: Added hashcat keyboard mapping de.hckmap
2018-11-15 22:29:03 +01:00
Jens Steube
5d5ac1c935
Prepare for on-the-fly keyboard layout substituations required to crack booting TrueCrypt/VeraCrypt volumes
2018-11-15 14:35:51 +01:00
Jens Steube
a4200ba167
Added hash-mode 18300 (Apple File System)
...
Fixes https://github.com/hashcat/hashcat/issues/1686
2018-11-12 11:37:01 +01:00
R. Yushaev
47bd838e25
Add VeraCrypt Streebog support
...
VeraCrypt added the possibility to use Streebog-512 as hashing algorithm
for the key derivation. This commit adds the necessary VeraCrypt kernels
as well as additional HMAC-Streebog kernels.
- Add hash-mode 13771: VeraCrypt PBKDF2-HMAC-Streebog-512 + XTS 512 bit
- Add hash-mode 13772: VeraCrypt PBKDF2-HMAC-Streebog-512 + XTS 1024 bit
- Add hash-mode 13773: VeraCrypt PBKDF2-HMAC-Streebog-512 + XTS 1536 bit
- Add hash-mode 11750: HMAC-Streebog-256 (key = $pass), big-endian
- Add hash-mode 11760: HMAC-Streebog-256 (key = $salt), big-endian
- Add hash-mode 11860: HMAC-Streebog-512 (key = $salt), big-endian
- Add test suite for hash-modes 11750, 11760 and 11860
- Improve pure Streebog kernels
2018-11-08 11:46:31 +01:00
Royce Williams
6053f473eb
trailing whitespace
2018-11-01 11:17:02 -08:00
R. Yushaev
a8eb611b1c
Add HMAC-Streebog-512 (pure kernels)
...
Implement HMAC based on GOST 34.11-2012 Streebog-512 as well as a test
case for it. Both the PyGOST + hmac python module and the VeraCrypt HMAC
for Streebog-512 were used as references. The kernels expect the digests
to be in big-endian order according to the RFC examples for Streebog.
Fix two bugs from commit 224315dd62
.
- Add hash-mode 11850: HMAC-Streebog-512 (key = $pass), big-endian
- Add test case for hash-mode 11850
- Bugfix for a3-pure Streebog kernels (modes 11700 and 11800)
- Rename a few Streebog constants in interface.h
2018-10-31 14:42:02 +01:00
Arseniy Sharoglazov
ee873da300
Added hash-modes 18200 (Kerberos 5 AS-REP etype 23)
2018-10-30 19:05:44 +03:00
R. Yushaev
224315dd62
Add pure kernels and tests for Streebog hashes
...
Complete Streebog support with pure kernels that allow for passwords
longer than 64 characters. Provide generic inc_hash_streebog files
for future Streebog-based hash modes (HMAC, PBKDF2, VeraCrypt).
Include streebog support in the test suite. For this, python module
PyGOST is needed. Also add clarification to hash mode description
stating that Streebog hashes are expected in big-endian byte order.
There are several implementations, including PyGOST, which default
to little-endian byte order, while the RFC examples are big-endian.
- Add pure kernels for hash-mode 11700 (Streebog-256)
- Add pure kernels for hash-mode 11800 (Streebog-512)
- Tests: Add hash-modes 11700 (Streebog-256) and 11800 (Streebog-512)
2018-10-29 10:33:30 +01:00
Jens Steube
5eca3f5316
Fix kernel names in -a 1 kernels of -m 7701 and 7801
2018-10-25 11:12:26 +02:00
Jens Steube
48cf3f722b
Merge pull request #1725 from unix-ninja/master
...
Cleanup unused register definitions
2018-10-23 09:43:06 +02:00
unix-ninja
6196e23069
Cleanup unused register definitions
2018-10-22 15:23:41 -04:00
Jens Steube
e2a9409413
Merge pull request #1710 from unix-ninja/master
...
Add support for TOTP (RFC 6238)
2018-10-22 20:49:31 +02:00
unix-ninja
8c5c225d8f
Optimize performance on NVIDIA GTX
2018-10-22 13:27:35 -04:00
Jens Steube
b0077860c7
Workaround some padding issues with host compiler and OpenCL JiT on 32 and 64 bit systems
2018-10-20 12:41:41 +02:00
Jens Steube
a4ac370496
Test fix for plain_t in 32 bit mode
2018-10-20 09:46:24 +02:00
Jens Steube
6d03da369b
Fix gid datatype in mark_hash()
2018-10-20 02:19:39 +02:00
unix-ninja
fddfd835d2
Support 64 bit timestamps properly
2018-10-19 15:35:52 -04:00
unix-ninja
7904b9ae05
Fix kernel types to align with style guide
2018-10-19 07:56:51 -04:00
Jens Steube
0a74f058ac
Synchronize salt_t datatypes in types.h and inc_types.cl
2018-10-19 10:20:13 +02:00