Jens Steube
58fa783095
Enhanced the auto-tune engine: when a kernel runs with a single thread and no accel, it should finish quickly (ideally under 1 ms). If it doesn't, the kernel is likely overloaded with code. If such a kernel also uses barriers (e.g., to load shared storage with multiple threads), high iteration counts cause unnecessary thread waiting. To address this, we now skip increasing the loop count if the runtime exceeds either 1/8 of the target time (based on the -w setting) or a hard-coded threshold of 4 ms.
...
Improved shared memory handling for -m 10700. Removed the hard-coded limit of 256 threads and now dynamically check the device's shared memory pool to adapt threads accordingly.
Implemented a feature request to display non-default session names early during startup.
Added a check for the number of registers required by a kernel (CUDA and HIP only). This allows us to estimate the max threads per block before entering the auto-tune engine and make pre-adjustments.
Fixed Metal command encoder argument to work with the new auto-tuner's extra kernel invocation.
Fixed incorrect host memory calculation logic during automatic kernel-accel reduction for scrypt-based algorithms. This ensures memory constraints are respected.
Improved several plugins by setting maximum loop counts and others using the OPTS_TYPE_NATIVE_THREADS option.
Fixed compilation on Apple platforms by excluding '#include <sys/sysinfo.h>'.
2025-06-25 22:10:29 +02:00
Jens Steube
15ada5124e
Further simplified the use of inc_hash_scrypt.cl without any speed regression, and updated all affected plugin kernels. Use m08900-pure.cl as a template.
...
Updated kernel declarations from "KERNEL_FQ void HC_ATTR_SEQ" to "KERNEL_FQ KERNEL_FA void". Please update your custom plugin kernels accordingly.
Added spilling size as a factor in calculating usable memory per device. This is based on undocumented variables and may not be 100% accurate, but it works well in practice.
Added a compiler hint to scrypt-based kernels indicating the guaranteed maximum thread count per kernel invocation.
Removed redundant kernel code 29800, as it is identical to 27700, and updated the plugin.
2025-06-21 17:41:26 +02:00
Jens Steube
4e0a728f8f
Add HC_ATTR_SEQ as a workaround, since HIP no longer
...
supports compiler option --gpu-max-threads-per-block
2025-06-12 09:17:02 +02:00
justpretending
b2f14f2f5d
Fix some typos
2023-07-27 23:11:55 +07:00
Gabriele Gristina
cf3ab8e2dc
Handle signed/unsigned PDF permission P value for all PDF hash-modes
2023-04-11 21:17:25 +02:00
Sein Coray
69807fb2fe
removed check which aborted on empty passwords for mode 10700
2023-01-09 09:00:26 +01:00
Jens Steube
f6537a2964
Use inline static on HIP for some hash-modes which benefit from it
2022-11-07 15:35:46 +01:00
philsmd
1b770fdea7
fix maximum password length in -O -m 10700 = PDF
2022-03-13 12:59:55 +01:00
Gabriele Gristina
b0c6738289
Rename STR() to M2S(), part 2
2022-02-07 13:31:22 +01:00
Jens Steube
be75e4b4ea
Rename STR() to M2S() to avoid future collisions and move from kernel source to command line parameter
2022-02-07 09:57:08 +01:00
Gabriele Gristina
9d36245d51
Kernels: Set the default Address Space Qualifier for any pointer, refactored / updated KERN_ATTR macros and rc4 cipher functions, in order to support Apple Metal runtime
2022-02-04 19:54:00 +01:00
Gabriele Gristina
2e4a136758
Refactored standard kernel includes in order to support Apple Metal runtime, updated backend, test units and status code
2022-01-18 22:52:14 +01:00
Jens Steube
0abdcb1ae5
Rename GID_MAX to GID_CNT to avoid naming conflict with existing macro
2022-01-04 22:57:26 +01:00
Jens Steube
668d2179cd
Kernels: Refactored standard kernel declaration to use a structure holding u32/u64 attributes to reduce the number of attributes
2022-01-04 18:07:18 +01:00
Jens Steube
cb7f99ef79
Renamed macro INLINE to HC_INLINE to avoid naming conflict with INLINE on MacOS
2021-12-24 16:40:43 +01:00
Jens Steube
89cd5bd78b
Remove inline static keyword in inc_vendor.h for HIP platform since it's the default setting with HIP 4.4
2021-11-02 08:12:13 +01:00
Jens Steube
aee8e559c4
PDF Kernel (10700): Improved performance on AMD GPU by using shared memory for the scratch buffer and disable inlining to save spilling
...
Inspired by https://github.com/reger-men/hashcat/blob/6.2.4/OpenCL/m10700-optimized.cl
2021-10-31 10:05:58 +01:00
Jens Steube
1d33b57144
PDF 1.7 Kernel: Improved performance on AMD GPU by using shared memory for the scratch buffer
...
Inspired by https://github.com/reger-men/hashcat/blob/6.2.4/OpenCL/m10700-optimized.cl
2021-10-30 20:16:45 +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
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
04d5e5a119
New Attack-Mode: Association Attack. Like JtR's single mode. Very early
...
stage. See hashcat Forum for detailed writeup.
2020-09-29 15:56:32 +02:00
Jens Steube
8ff8c5d536
Add LOCAL_VK to make use of __shared__
2019-05-07 09:01:32 +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
jsteube
7c6970dbdd
Remove hard-coded static keyword from OpenCL kernels
2019-04-13 18:46:19 +02:00
jsteube
1c45de4b63
Remove more old GCC diagnostic options
2019-03-31 19:41:28 +02:00
Jens Steube
0fb3b3c83e
Declare internal functions in OpenCL kernels as static
2019-03-26 11:03:25 +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
c9d60c079f
Prepare OpenCL kernels for non-static compilation
2019-03-22 15:16:25 +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
Jens Steube
09b8bfc3a5
Fix datatypes in -m 10700
2019-03-08 13:13:01 +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
e1fe3e755b
Optimize some switch_buffer_* functions for generic OpenCL devices (CPU, various OSX, ...)
2019-03-01 14:49: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
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
2a6444c05a
Give the compiler a hint for automatic optimizations based on password length
2018-11-20 15:26:46 +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
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
jsteube
68bff94980
Workaround rocm OpenCL runtime bug when copy data from constant to local memory
2018-08-12 18:04:33 +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
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
3e08750900
OpenCL Kernels: Add general function declaration keyword (static inline)
2018-02-06 19:12:24 +01: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
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
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
4b2f3011e9
Mix in pure kernel functions to PDF 1.7 Level 8 (Acrobat 10 - 11)
2017-08-10 10:43:30 +02:00
jsteube
fbea72ebd6
Renamed default kernels to optimized kernels
...
Renamed pure kernels to default kernels
Replaced long option --length-limit-disable with --optimized-kernel-enable
Replaced short option -L with -O
Set --optimized-kernel-enable to unset by default
2017-07-18 13:23:42 +02:00
