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 |
|
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 |
|
Jens Steube
|
bd92589af1
|
Optimize BLAKE2B_ROUND() 64 bit rotates
|
2021-07-17 18:18:22 +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 |
|
Jens Steube
|
3becb253d3
|
Enable vector datatypes for HIP
|
2021-07-17 18:01:01 +02: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 |
|
Gabriele Gristina
|
0b19c4a688
|
fix m26700_comp, cleanup changes.txt
|
2021-07-17 14:44:36 +02: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
|
5b9823633b
|
Add missing entry about HIP to docs/changes.txt
|
2021-07-16 22:37:29 +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
|
b25b7c2506
|
Add hashcat.hctune entries for scrypt based hash-modes for AMD MI100
|
2021-07-16 17:55:12 +02:00 |
|
Michael Kruger
|
39e78702e0
|
Move gid check below SHM stuff to imitate m06232
|
2021-07-16 16:01:57 +02:00 |
|
Michael Kruger
|
a6ae33b1cd
|
Changed the test cases to expect the HEX
|
2021-07-16 15:38:42 +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
|
0f055d93ba
|
fix 26700 struct name
|
2021-07-15 20:59:54 +02:00 |
|
Gabriele Gristina
|
eaedfb3f8a
|
fix 26700
|
2021-07-15 19:58:11 +02:00 |
|
Cablethief
|
32b4a2cfd5
|
Modded test, same issue about requireing a HEX input or HEX[] input
|
2021-07-15 16:50:26 +02:00 |
|
Cablethief
|
431f6ec80e
|
remove fast hash 27000
|
2021-07-15 16:43:48 +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 |
|
Cablethief
|
c2ceba4c9d
|
Modded test, however unsure of how to specify only HEX input or not
|
2021-07-15 15:56:39 +02:00 |
|
Jens Steube
|
372ca6609f
|
Merge pull request #2881 from matrix/ext_sysfs_cpu_fix
Fix bug on hm_SYSFS_CPU_get_syspath_hwmon()
|
2021-07-15 13:23:48 +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 |
|