Commit Graph

7670 Commits

Author SHA1 Message Date
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
Jens Steube 640d95a00f Vendor Detection: Add "Intel" as a valid vendor name for GPU on macOS 2021-07-24 16:01:30 +02:00
Jens Steube 95af9964bf Merge pull request #2896 from matrix/snmpv3_md5
Added hash-mode: SNMPv3 HMAC-MD5-96
2021-07-24 15:13:28 +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 edab99cdfa Update SCRYPT settings in hashcat.hctune for MI100 and RX6900XT 2021-07-22 18:57:36 +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
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 a09efb2634 SCRYPT Kernels: Add more optimized values for some new NV/AMD GPUs 2021-07-22 11:46:11 +02:00
Jens Steube 0d64db07d5 Merge pull request #2892 from Xanadrel/patch-1
Add hashcat.hctune entries for Nvidia RTX 3060Ti & 3070 for scrypt based hash-modes
2021-07-22 11:38:41 +02:00
Xanadrel 1d7880fa30 Add hashcat.hctune entries for Nvidia RTX 3060Ti & 3070 for scrypt based hash-modes 2021-07-21 23:24:19 +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
Jens Steube 8ba907125a Merge pull request #2888 from jtojanen/master
Calculate kernel dynamic memory size based on CU_DEVICE_ATTRIBUTE_MAX_SHARED_MEMORY_PER_BLOCK_OPTIN
2021-07-21 13:28:12 +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 db7cb16a3b using sha1_hmac_update_global 2021-07-21 10:19:05 +02: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 41156bb0a5 Improve PIM fix for use on macOS 2021-07-20 15:10:49 +02: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
Jens Steube 777512e862 Fixed false negative in all VeraCrypt hash-modes if both conditions are met: 1. use CPU for cracking and 2. PIM range was used 2021-07-20 11:31:31 +02: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 92dc57df28 Revert adding make_u32x() for constants to crypto primitives, fails on OpenCL 2021-07-20 10:34:34 +02:00
Jens Steube bdb7999f07 Switch HIP vector datatypes to OpenCL like ext_vector_type() 2021-07-19 20:24:30 +02:00
Jens Steube 9421e1f89d Fixed datatype in function sha384_hmac_init_vector_128() that could come into effect if vector datatype was manually set 2021-07-19 15:58:38 +02: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 0d8b4b74ad More CUDA special backports to HIP 2021-07-18 22:56:22 +02:00
Jens Steube 257098a301 Get rid of hip/hip_runtime.h dependancy 2021-07-18 21:14:45 +02:00