1
0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-26 09:58:16 +00:00
Commit Graph

1939 Commits

Author SHA1 Message Date
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