arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-23 23:19:09 +00:00
Code Issues Releases Wiki Activity
7,908 Commits 1 Branch 27 Tags 209 MiB
63cc905bbc
Commit Graph

4177 Commits

Author SHA1 Message Date
Jens Steube
af5d346244
Merge pull request #2894 from jtojanen/master 
Allow async execution
 2021-08-02 10:05:20 +02:00
therealartifex
483a62cb7c Add test module, update name of hash mode 27100 2021-08-01 21:31:06 -04:00
therealartifex
dc662c354e Update hash-mode number 2021-08-01 21:09:56 -04:00
Jens Steube
20a7b9f992 Tuning-Database: Add new module function module_extra_tuningdb_block() to extend hashcat.hctune content from a plugin 
See src/modules/module_08900.c as an example
 2021-08-01 16:25:37 +02:00
Jukka Ojanen
3d7ce7162b Merge branch 'master' of https://github.com/hashcat/hashcat 2021-08-01 15:36:52 +03:00
Jukka Ojanen
62a06f735f kernel_memset32(): assume offset and size are bytes, not elements 2021-08-01 15:19:15 +03:00
Jukka Ojanen
fdbfae9a28 Modify OpenCL clEnqueueFillBuffer() workaround 2021-08-01 14:47:08 +03:00
Jukka Ojanen
dbe2bad098 Maintain code style 2021-08-01 14:26:33 +03:00
Jukka Ojanen
1ee222d43f Replace free() with hcfree() 2021-08-01 14:21:18 +03:00
Jukka Ojanen
e352a79a05 Extend context in gidd_to_pw_t() 2021-08-01 14:15:03 +03:00
Jens Steube
9e077575c6
Merge branch 'master' into fix_24700_bof 2021-08-01 10:31:37 +02:00
Jens Steube
a440a4dee5
Merge branch 'master' into fix_25500_bof 2021-08-01 10:27:23 +02:00
Jens Steube
6bcbc218d6 Fixed out-of-boundary read in input_tokenizer() if the signatures in the hash line is longer than the constant signature in the plugin 2021-08-01 10:21:21 +02:00
Jens Steube
f4e52ca533 Add new rule function '3' to switch the case of the first letter after occurrence N of char X 2021-08-01 00:04:10 +02:00
Jens Steube
d4997d1255 Added support for auto-tuning --kernel-threads (-T) on startup 2021-07-31 16:49:39 +02:00
Gabriele Gristina
89234a5c7c Fixed buffer overflow in Stargazer Stellar Wallet XLM module 2021-07-31 16:18:29 +02:00
Gabriele Gristina
4869e2a9d1 Fixed buffer overflow in Stuffit5 module 2021-07-31 15:52:44 +02:00
Gabriele Gristina
a847a4c84b Fixed multiple buffer overflow in DPAPI masterkey file v2 module 2021-07-31 14:38:14 +02:00
Gabriele Gristina
ca896f89bc Fixed multiple buffer overflow in DPAPI masterkey file v1 module 2021-07-31 14:34:36 +02:00
therealartifex
6d4dbffda1 Moved the addition of static dashes to the kernel 2021-07-30 17:13:01 -04:00
therealartifex
1c9add307a
Merge branch 'hashcat:master' into master 2021-07-30 15:07:50 -04:00
Jens Steube
bb1ead3b01
Merge branch 'master' into snmpv3_sha256 2021-07-30 18:57:44 +02:00
Jens Steube
8801855b14
Merge branch 'master' into snmpv3_sha224 2021-07-30 18:56:20 +02:00
therealartifex
73104fafea Update based on recommendations and examination of mode 19500 2021-07-30 10:40:21 -04:00
therealartifex
33ce4f9720 Add new hash-mode 28200 (SHA1DASH), used in Zynga, 8tracks, and avvo 2021-07-30 10:39:45 -04:00
philsmd
92dfa7e137 Increase buffer sizes for -m 11600 = 7-Zip 2021-07-30 15:29:47 +02:00
philsmd
51adbbbe53 fix buffer sizes for -m 13600 = WinZip 2021-07-30 15:23:33 +02:00
Jukka Ojanen
e154f9e781 Few cleanups 2021-07-30 12:24:21 +03:00
Jukka Ojanen
c3195d0603 Merge branch 'master' of https://github.com/hashcat/hashcat 2021-07-30 11:34:25 +03:00
Jukka Ojanen
d721deb17a Move variables higher to avoid async execution stack problems 2021-07-29 15:23:24 +03:00
Jukka Ojanen
1064cce08c Synchronize before hooks 2021-07-29 15:22:25 +03:00
Jukka Ojanen
81c2ec3caf Small cleanup in gidd_to_pw_t() 2021-07-29 14:48:36 +03:00
Jukka Ojanen
8d51fca192 No need to synchronize after run_kernel() 2021-07-29 14:40:45 +03:00
Jukka Ojanen
7a8065d090 Do not call clWaitForEvents() after spin damper when we know that event status is CL_COMPLETE 2021-07-29 14:39:11 +03:00
Jukka Ojanen
c48e6a25a8 Enqueue several commands before clFlush() 2021-07-29 14:28:01 +03:00
Jens Steube
03ed06849a
Merge branch 'HIP44' into master 2021-07-29 11:00:28 +02:00
Jens Steube
d38d40c8ba Unlock all GPU threads for AMD GPUs if WaveFront size is 32 (basically new models) 
Add new hash-modes to tools/benchmark_deep.pl
Fix MINGW issue on 64 bit constant in refactored kernel-accel limiting section
 2021-07-29 10:49:44 +02:00
Jukka Ojanen
9ed231c99c Add comment to blocking OpenCL calls 2021-07-29 00:07:10 +03:00
Jukka Ojanen
e889cf98a9 Few cleanups in autotune 2021-07-28 23:50:58 +03:00
Jens Steube
a4299b74af Memory Management: Refactored the code responsible for limiting kernel accel in order to avoid out of -host- memory situations 2021-07-28 20:50:05 +02:00
Jukka Ojanen
d7cc8d7cd2 Revert due to module_03200.c, module_25600.c and module_25800.c using device_param->kernel_dynamic_local_mem_size_memset 2021-07-28 19:49:02 +03:00
Jukka Ojanen
8b590f651b Remove unused gpu_memset and its references 2021-07-28 19:26:30 +03:00
Jens Steube
9c0a37accf Update driver requirement 2021-07-28 07:56:34 +02:00
Jens Steube
ae39fa0979
Merge pull request #2908 from hashcat/master 
Backports
 2021-07-28 07:55:11 +02:00
Jens Steube
25f1c12e3c SCRYPT Kernels: Add more optimized values for some new NV/AMD GPUs and new semi-automated derivation process description 
Blowfish Kernels: Backport optimizations reducing bank conflicts from bcrypt to Password Safe v2 and Open Document Format (ODF) 1.1
 2021-07-28 07:51:27 +02:00
Jukka Ojanen
5dbe4958b8 Remove digests_shown_tmp buffer, used in check_cracked() to zero out hashes 2021-07-27 19:01:13 +03:00
Jukka Ojanen
cdf27a1cb3 Implement async run_cuda_kernel_memset() and run_cuda_kernel_memset32() 2021-07-27 18:56:59 +03:00
Jukka Ojanen
e8be7028cd Handle errors inside check_cracked() 2021-07-27 16:28:07 +03:00
Jukka Ojanen
be33ccaa6c check_hash() needs to return value to indicate possible error 2021-07-27 14:36:48 +03:00
Jens Steube
532a154542 ADL: Updated support for AMD Display Library to 15.0, updated datatypes and added support for OverDrive 7 and 8 based GPUs 2021-07-27 12:02:27 +02:00
Jukka Ojanen
ae44a3022b Use fast event to synchronize selftest() 2021-07-27 13:00:39 +03:00
Jukka Ojanen
d7de3550b1 Add fast event for CUDA and HIP, and use it to synchronize check_hash() 2021-07-27 12:57:26 +03:00
Jens Steube
f6de3e61e0
Merge pull request #2906 from hashcat/master 
Backports
 2021-07-27 09:48:56 +02:00
Jens Steube
e5ac568afb
Merge pull request #2903 from matrix/snmpv3_mod 
fix snmpv3 md5/sha1, tested with real hashes
 2021-07-27 09:47:47 +02:00
Jens Steube
e627288c04
Merge pull request #2897 from matrix/snmpv3_md5_sha1 
Added hash-mode: SNMPv3 HMAC-MD5-96/HMAC-SHA1-96
 2021-07-27 09:47:11 +02:00
Jens Steube
fd2cb59d26 AMD GPUs: On Apple OpenCL platform, we ask for the preferred kernel thread size rather than hard-coding 32 
ECC secp256k1: Removed the inline assembly code for AMD GPUs because the latest JIT compilers optimize it with the same efficiency
 2021-07-27 09:37:31 +02:00
Gabriele Gristina
357c23d7a1 Added hash-mode: SNMPv3 HMAC-SHA384-256 2021-07-27 05:17:26 +02:00
Gabriele Gristina
15b90c953e Added hash-mode: SNMPv3 HMAC-SHA512-384 2021-07-27 05:14:09 +02:00
Gabriele Gristina
d711c593de fix to crack real hashes 2021-07-26 23:00:33 +02:00
Gabriele Gristina
d62fa617fb fix to crack real hashes 2021-07-26 22:45:02 +02:00
Gabriele Gristina
c20ff01c39 using shared buffer between md5 and sha1 SNMPV3_TMP_ELEMS_OPT, fix to crack real hashes 2021-07-26 22:25:15 +02:00
Jukka Ojanen
73b8cda176 Avoid unnecessary memcpy in check_hash() 2021-07-26 22:25:50 +03:00
Jukka Ojanen
43a3622a77 Add necessary synchronization to check_hash() and check_cracked() 2021-07-26 20:57:24 +03:00
Gabriele Gristina
e15fe3461e add missing pw_min() to snmpv3 md5/sha1 modules 2021-07-26 19:38:53 +02:00
Gabriele Gristina
45fce5d3a3 fix snmpv3 md5/sha1, tested with real hashes 2021-07-26 19:04:30 +02:00
Jukka Ojanen
a642f7b233 Remove synchronous GPU memory copy functions 2021-07-26 15:36:42 +03:00
Jukka Ojanen
507d40fff6 Allow async execution of check_hash() and check_cracked() 2021-07-26 15:35:59 +03:00
Jukka Ojanen
de5200cffc Allow async execution of backend 2021-07-26 15:17:25 +03:00
Jens Steube
7f419c68af Blowfish Kernels: Backport optimizations reducing bank conflicts from bcrypt to Password Safe v2 and Open Document Format (ODF) 1.1 2021-07-26 13:38:39 +02:00
Jens Steube
b66c093c17
Merge pull request #2902 from hashcat/master 
Backports
 2021-07-26 08:18:35 +02:00
Jens Steube
1e3bd2c8a0 AMD GPUs: Add inline assembly code for md5crypt/sha256crypt, PDF 1.7, 7-Zip, RAR3, Samsung Android and Windows Phone 8+ 2021-07-26 07:59:12 +02:00
Jens Steube
b53691c8f5 ADL: Updated support for AMD Display Library to 14.0, updated datatypes and added support for OverDrive 7 and 8 based GPUs 2021-07-26 07:48:56 +02:00
Jukka Ojanen
a86b44a750 Allow async execution of autotune() and prepare try_run() for async kernel execution 2021-07-26 01:51:19 +03:00
Jukka Ojanen
79404b1ff1 Allow async execution of selftest() 2021-07-25 22:19:50 +03:00
Jens Steube
72e307fbce
Merge pull request #2901 from hashcat/master 
Backport changes
 2021-07-25 10:57:39 +02:00
Jens Steube
83badbeaf1 Backport native threads strategy from Kerberos to Office and PDF 2021-07-25 10:56:29 +02:00
Jens Steube
959a232828
Merge pull request #2885 from neheb/charfixes 
const and char fixes
 2021-07-25 10:36:40 +02:00
Jens Steube
84a4058edf
Merge pull request #2900 from hashcat/master 
Backport changes
 2021-07-25 10:34:05 +02:00
Jens Steube
44dbedd21a Use improved native threads derivation for RC4 based hash-modes 2021-07-25 10:32:58 +02:00
Gabriele Gristina
58bb2b69b4 Added hash-mode: SNMPv3 HMAC-SHA256-192 2021-07-24 18:57:25 +02:00
Gabriele Gristina
7683ca19a3 Added hash-mode: SNMPv3 HMAC-SHA224-128 2021-07-24 18:36:10 +02:00
Gabriele Gristina
7c8d7f7831 Added hash-mode: SNMPv3 HMAC-MD5-96/HMAC-SHA1-96 2021-07-24 17:03:50 +02:00
Jens Steube
640d95a00f Vendor Detection: Add "Intel" as a valid vendor name for GPU on macOS 2021-07-24 16:01:30 +02:00
Gabriele Gristina
4f10377703 Added hash-mode: SNMPv3 HMAC-MD5-96 2021-07-24 13:56:49 +02:00
Jens Steube
3a31b669b5 Optimize and clean up -m 25200 code. 
Speed on CPU: 1521 -> 2066 (35% faster)
Speed on GPU: 16610 -> 157754 (9.4 times faster)
Works for all password lengths
 2021-07-24 11:58:36 +02:00
Jens Steube
5e0f7ecf00
Merge pull request #2889 from matrix/snmpv3_sha1 
Added hash-mode: SNMPv3 HMAC-SHA1-96
 2021-07-24 11:56:46 +02:00
Gabriele Gristina
4fb44f826c updated module_25200.c as requested in the review 2021-07-24 10:57:06 +02:00
Jens Steube
3f6c5a0042 Update module_unstable_warning() for -m 172xx on HIP 2021-07-23 21:09:55 +02:00
Jens Steube
5ffcaa980d HIP Backend: Added support to support HIP 4.4 and later, but added check to rule out older versions because they are incompatible 2021-07-23 16:04:34 +02:00
Rosen Penev
adaf3f293b make const char pointers actually const 
const char* is a non const pointer that points to const data. Add
missing const.

Signed-off-by: Rosen Penev <rosenp@gmail.com>
 2021-07-22 18:13:46 -07:00
Rosen Penev
a9ceb6377e add missing const 
Fixes -Wwrite-strings issue.

Signed-off-by: Rosen Penev <rosenp@gmail.com>
 2021-07-22 18:11:11 -07:00
Rosen Penev
14f5a26027 use const char for fopen mode 
Fixes -Wwrite-strings warnings.

Signed-off-by: Rosen Penev <rosenp@gmail.com>
 2021-07-22 18:11:11 -07:00
Gabriele Gristina
b00e3e5e59 add missing hex_encode() in module_hash_encode() 2021-07-22 19:53:41 +02:00
Gabriele Gristina
6d83a69f7d using hex_encode/hex_decode + small fixes 2021-07-22 19:10:34 +02:00
Jens Steube
f9e74045b5 Commandline: Throw an error if separator character given by the user with -p option is not exactly 1 byte 2021-07-22 18:54:02 +02:00
Jukka Ojanen
9f9333f2ef Allow async execution of run_opencl_kernel_bzero(), run_hip_kernel_bzero() and run_opencl_kernel_bzero() 2021-07-22 15:18:10 +03:00
Jukka Ojanen
fafd24237d Define HC_ALIGN macro to control data alignment and use common constant zero buffer in run_cuda_kernel_bzero(), run_hip_kernel_bzero() and run_opencl_kernel_bzero(). 2021-07-22 14:24:03 +03:00
Jukka Ojanen
2c2988518d Remove all calls to clFlush() 2021-07-22 13:59:19 +03:00
Jens Steube
5024865d87 Kernel Threads: Use warp size / wavefront size query instead of hardcoded values as base for kernel threads 
Kernel Cache: Add kernel threads into hash computation which is later used in the kernel cache filename
Remove some unused function symbol lookups in HIP library
 2021-07-22 11:46:47 +02:00
Jens Steube
e83611742d Brain Session: Adds hashconfig specific opti_type and opts_type parameters to hashcat session computation to cover features like -O and -M 2021-07-21 15:47:26 +02:00
Jens Steube
c990e252d3 Added option --multiply-accel-disable (short: -M) to disable multiply the kernel-accel with the multiprocessor count automatism 2021-07-21 15:47:05 +02:00
Jens Steube
a7a899e5a4 Backport changes from #2888 to HIP backend 2021-07-21 14:01:28 +02:00
Jukka Ojanen
cb923d6e46 Replace CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK with CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN 2021-07-21 13:42:22 +03:00
Gabriele Gristina
7ea5da51a8 Added hash-mode: SNMPv3 HMAC-SHA1-96 2021-07-20 22:24:35 +02:00
Jukka Ojanen
55a475cb4b
Merge branch 'hashcat:master' into master 2021-07-20 21:27:51 +03:00
Jukka Ojanen
d23f2d6c2f Calculation kernel dynamic memory size based on CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK 2021-07-20 21:27:29 +03:00
Jens Steube
a026171fb3
Merge pull request #2879 from jtojanen/master 
Implement gpu_bzero
 2021-07-20 13:00:37 +02:00
Jens Steube
7c5fe5e233 Slow Kernels: Set some of the slowest kernels to OPTS_TYPE_MP_MULTI_DISABLE 2021-07-20 12:47:40 +02:00
Jukka Ojanen
8674e23d79 Add async HIP memcpy functions: hc_hipMemcpyDtoDAsync(), hc_hipMemcpyDtoHAsync() and hc_hipMemcpyHtoDAsync(). Implement partially async HIP memset and bzero kernels. 2021-07-20 12:47:10 +03:00
Jukka Ojanen
4263cafdcf Add async CUDA memcpy functions: hc_cuMemcpyDtoDAsync(), hc_cuMemcpyDtoHAsync() and hc_cuMemcpyHtoDAsync(). Implement partially async CUDA memset and bzero kernels. 2021-07-20 12:23:39 +03:00
Jukka Ojanen
f07ff6f03d Use kernel_bzero instead of kernel_memset when value is zero 2021-07-20 12:04:12 +03:00
Jukka Ojanen
4c0f6dd263
Merge branch 'hashcat:master' into master 2021-07-20 12:00:41 +03:00
Jukka Ojanen
ea5180ac46 Include missing bzero kernel parameters 2021-07-20 11:59:44 +03:00
Jens Steube
7d43b6e1eb Clear tmps memory region after autotune use 2021-07-19 11:59:53 +02:00
Jens Steube
0504498e86
Merge pull request #2886 from pellekuiters/issue-2876 
Auto-tuning outside kernel can cause unpredictable behavior for tmp_t structs on CUDA
 2021-07-19 11:52:07 +02:00
pelle
ac3ccfcc51 Run init() and loop_prepare() during autotune to prevent possible corruptions during loop(). 2021-07-19 11:07:24 +02:00
Jens Steube
257098a301 Get rid of hip/hip_runtime.h dependancy 2021-07-18 21:14:45 +02:00
Gabriele Gristina
977ceba046 Added hash-modes: VirtualBox (PBKDF2-HMAC-SHA256 & AES-128-XTS) and VirtualBox (PBKDF2-HMAC-SHA256 & AES-256-XTS) 2021-07-17 18:38:04 +02:00
Jukka Ojanen
72a418932c HIP: Use kernel_bzero instead of kernel_memset when value is zero 2021-07-17 19:12:09 +03:00
Jukka Ojanen
a2a1d04bcf Implement gpu_bzero 2021-07-17 19:00:10 +03:00
Jukka Ojanen
8066a47ac5
Merge branch 'hashcat:master' into master 2021-07-17 16:21:08 +03:00
Jens Steube
261e0f42ab Unset PRODUCTION option 2021-07-16 23:03:11 +02:00
Jens Steube
84082a952e hashcat 6.2.3 2021-07-16 22:41:52 +02:00
Jens Steube
141b59b80e
Merge pull request #2883 from hashcat/HIP 
Merge HIP branch into master
 2021-07-16 19:54:00 +02:00
Jens Steube
45e65dd05a Backport more ROCm based optimizations to HIP 2021-07-15 23:34:27 +02:00
Jens Steube
2e929e692e Backport update module_unstable_warning() for -m 21800 on HIP 2021-07-15 21:26:07 +02:00
Gabriele Gristina
eaedfb3f8a fix 26700 2021-07-15 19:58:11 +02:00
Jens Steube
4730cf6e79 WinZip Kernel: Increase supported data length from 8k to 16mb 2021-07-15 16:42:29 +02:00
Cablethief
02a9f3471b 27000 slow hash, same doubts as the first, but now with the shared mem concerns 2021-07-15 16:42:18 +02:00
Jens Steube
d130cc66b3 Optimize ISA code on HIP for V_ALIGNBIT_B32 using a different template for inline assembly 2021-07-15 09:57:41 +02:00
Gabriele Gristina
50aeaa299d Added hash-mode: VMware VMX (PBKDF2-HMAC-SHA1 + AES-256-CBC) 2021-07-14 19:31:21 +02:00
Gabriele Gristina
80a671eab7 Fix bug on hm_SYSFS_CPU_get_syspath_hwmon() 2021-07-14 19:21:21 +02:00
Jens Steube
cf512faa53 Update large switch() cases in inc_common.cl and some inline assembly common functions for devices managed with HIP backend 2021-07-14 17:06:20 +02:00
Jens Steube
11295e4679 Fix missing OPTI_TYPE_USES_BITS_64 in several modules 2021-07-14 17:01:46 +02:00
Jukka Ojanen
43b5a138d4 Use kernel_bzero instead of kernel_memset when value is zero 2021-07-14 17:13:39 +03:00
Cablethief
a38b5b3e25 27100 Slow hash, have doubts on how to correctly set iterations 2021-07-14 14:58:26 +02:00
Jens Steube
a82fce2c8f Fixed invalid ETA if --limit was specified, several salts are in a hash list and some of them were found in a potfile 2021-07-14 10:26:57 +02:00
Jens Steube
9c134833a6 Update module_unstable_warning() for -m 19600 and -m 19700 2021-07-14 08:26:12 +02:00
Jens Steube
f3bd936971 Add hardware monitor mapping for HIP devices 2021-07-14 08:23:39 +02:00
Jens Steube
84fcea9592 Fix unusued variable warning 2021-07-14 08:17:55 +02:00
Jens Steube
4b68a8fb4f
Merge pull request #2878 from hashcat/master 
Backports
 2021-07-13 22:26:28 +02:00
Jens Steube
7faf6859d6 Backport hand-optimized compiler settings in modules from ROCM to HIP 
Backport DECLSPEC settings from ROCM to HIP
 2021-07-13 20:45:01 +02:00
Jens Steube
74fca7bac1
Merge pull request #2875 from jtojanen/master 
Fix leaks in dictstat_destroy
 2021-07-13 13:28:37 +02:00
Jens Steube
219bed457f Fix use of --gpu-max-threads-per-block 2021-07-12 14:02:43 +02:00
Jukka Ojanen
6b5af5e849 Fix leaks in dictstat_destroy 2021-07-12 14:33:24 +03:00
Jens Steube
674ca7d88f Add GPU threads to kernel cache checksum because it has an influence on HIP offline compile options 
Add V_ALIGNBIT_B32 inline assembly wrapper because HIP does not provide amd_bitalign()
 2021-07-12 11:27:05 +02:00
Jens Steube
23c3c178bf Limit max threads per block to 64 to enable offline compiler to make better use if available registers 
Fix double free() for hip_event1/hip_event2 and hip_stream causes segfault
Replace hc_cuCtxSetCurrent() with hc_cuCtxPushCurrent() in order to align changes with HIP
Add vector datatype operators (if we decide to use them - currently unused)
 2021-07-12 09:28:26 +02:00
Jens Steube
20f7febd4c Workaround too intensive optimization in -m 2000 using HIPRTC 2021-07-11 15:54:13 +02:00
Jens Steube
fde4770219 Add misssing HIP backend information in -I output 2021-07-11 14:47:48 +02:00
Jens Steube
ca3beacd93 Disable dynamic shared memory on HIP, because hipFuncSetAttribute() maps to cudaFuncSetAttribute() and not to cuFuncSetAttribute() 2021-07-11 14:30:49 +02:00