arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-27 00:48:13 +00:00
Code Issues Releases Wiki Activity
7,679 Commits 1 Branch 27 Tags 209 MiB
25f1c12e3c
Commit Graph

2332 Commits

Author SHA1 Message Date
Gabriele Gristina
af622f6df5 Added hash-mode: BitShares v0.x - sha512(sha512(pass)) 2019-07-26 23:47:32 +02:00
Gabriele Gristina
bbee1890cd Added hash-mode: sha1(md5(pass).salt) 2019-07-26 01:49:35 +02:00
Gabriele Gristina
eedceb698f Added hash-mode: md5(sha1().md5().sha1()) 2019-07-25 23:56:20 +02:00
Gabriele Gristina
c166242996 Add AuthMe - sha256(sha256().) kernel module 2019-07-23 19:23:06 +02:00
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
unix-ninja
0e5704c77e Disable NEW_SIMD_CODE for 18100 (it's not compatible) 2018-10-18 15:58:21 -04:00
unix-ninja
b657c75583 Explicity cast otp_offset 2018-10-18 09:36:57 -04:00
unix-ninja
3869ce9246 More coding style fixes 2018-10-18 08:55:55 -04:00
unix-ninja
24ab7cae2a Add a1 kernel for mode 18100 2018-10-17 16:47:58 -04:00
unix-ninja
db4ec8ed2c Fix formatting to comply with hashcat coding guidelines 2018-10-17 16:34:34 -04:00
unix-ninja
1d43540fc4 Simplify alignment masks for mode 18100 2018-10-17 11:03:20 -04:00
unix-ninja
b29b7b8188 Convert arithmetic ops to logical ops in byte alignment 2018-10-17 08:54:52 -04:00
unix-ninja
3c3b05d1e5 Resolve conflicts 2018-10-16 15:48:20 -04:00
unix-ninja
6cda8f7077 Change TOTP index from 17300 to 18100 2018-10-16 15:33:09 -04:00
unix-ninja
73aae1a734 Add a0 kernel for TOTP 2018-10-16 15:07:41 -04:00
unix-ninja
2249ab4c13 Cleanup debug code 2018-10-16 15:05:21 -04:00
unix-ninja
977b560bb4 Add support for TOTP (RFC 6238) 2018-10-16 15:05:14 -04:00
R. Yushaev
5c87720acc Add SHA3 and Keccak 
The previous hash-mode 5000 covered Keccak-256 only. FIPS changed one
padding byte while adopting Keccak as the SHA3 standard, which gives us
different digests. Now we have separate kernels for SHA3 and Keccak.

 - Added hash-mode 17300 = SHA3-224
 - Added hash-mode 17400 = SHA3-256
 - Added hash-mode 17500 = SHA3-384
 - Added hash-mode 17600 = SHA3-512
 - Added hash-mode 17700 = Keccak-224
 - Added hash-mode 17800 = Keccak-256
 - Added hash-mode 17900 = Keccak-384
 - Added hash-mode 18000 = Keccak-512
 - Removed hash-mode 5000 = SHA-3 (Keccak)
 2018-10-15 16:06:31 +02:00
Michael Sprecher
1892b842d7
Increased the maximum size of edata2 in Kerberos 5 TGS-REP etype 23 2018-09-12 12:25:02 +02:00
jsteube
6e1aec0563 Fix kernel name in 16801 kernel source 2018-09-02 12:43:53 +02:00
Jens Steube
466ea8eaba Fixed detection of unique ESSID in WPA-PMKID-* parser 2018-08-31 15:47:48 +02:00
Michael Sprecher
5536ab9917
Getting rid of OPTS_TYPE_HASH_COPY for Ansible Vault 2018-08-15 23:32:58 +02:00
jsteube
a5746548e8 Allow use of hash-mode 7900, 10700 and 13731 on AMD devices after workaround 2018-08-13 13:41:43 +02:00
jsteube
6469357c74 Remove SCR_TYPE macro from OpenCL code 
Disable REAL_SHM access to AMD platform devices
 2018-08-13 12:10:03 +02:00
jsteube
68bff94980 Workaround rocm OpenCL runtime bug when copy data from constant to local memory 2018-08-12 18:04:33 +02:00
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