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
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
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
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
Jens Steube
4730cf6e79
WinZip Kernel: Increase supported data length from 8k to 16mb
2021-07-15 16:42:29 +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
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
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
4e773f32b2
Fix variable declaration block level in -m 19600 and -m 19700
2021-07-13 22:22:53 +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
1ab56afb9e
Merge pull request #2874 from hashcat/master
...
Backports from master
2021-07-12 08:57:50 +02:00
Jens Steube
5d05c95a2f
Fix address space of crc32tab[] in -m 172xx
2021-07-12 08:56:10 +02:00