1
0
mirror of https://github.com/hashcat/hashcat.git synced 2025-07-04 05:42:35 +00:00
Commit Graph

9769 Commits

Author SHA1 Message Date
Jens Steube
c59d3b8f34
Merge pull request #4271 from matrix/backend_info_machine_readable
Backend Info: Added --machine-readable format
2025-06-29 14:58:49 +02:00
Jens Steube
0c2ed0d199 Update plugins that benefit from an artificially limited register count (NVIDIA).
Update default hash settings to 64MiB:3:4 for Argon2 in -m 70000, following RFC 9106 recommendations.
Add option OPTS_TYPE_THREAD_MULTI_DISABLE: allows plugin developers to disable scaling the password candidate batch size based on device thread count. This can be useful for super slow hash algorithms that utilize threads differently, e.g., when the algorithm allows parallelization. Note: thread count for the device can still be set normally.
Add options OPTI_TYPE_SLOW_HASH_DIMY_INIT/LOOP/COMP: enable 2D launches for slow hash init/loop/comp kernel with dimensions X and Y. The Y value must be set via salt->salt_dimy attribute.
Change autotune kernel-loops start value to the lowest multiple of the target hash iteration count, if kernel_loops_min permits.
Fixed a bug in autotune where kernel_threads_max was not respected during initial init and loop-prepare kernel runs.
2025-06-29 14:39:14 +02:00
Jens Steube
237d784742
Merge pull request #4272 from matrix/regsPerBlock_workaround
workaround for HIP bug and avoiding a potential same bug on CUDA
2025-06-28 23:03:13 +02:00
Gabriele Gristina
c275c35ced
workaround for HIP bug and avoiding a potential same bug on CUDA 2025-06-28 22:54:36 +02:00
Gabriele Gristina
92b2f996b2
update docs/changes.txt 2025-06-28 10:45:24 +02:00
Gabriele Gristina
45b8672270
add --machine-readable format to --backend-info 2025-06-28 10:42:47 +02:00
Jens Steube
7fff4c929a Fixed a division by zero bug in fast hashes caused by hashes->st_salts_buf->salt_iter not being used. 2025-06-28 07:04:44 +02:00
Jens Steube
d5050d1f32
Merge pull request #4218 from matrix/ripemd320
Added hash-modes: RIPEMD-320, HMAC-RIPEMD320 (key = $pass), HMAC-RIPEMD320 (key = $salt)
2025-06-27 22:05:50 +02:00
Jens Steube
92ed2c6524
Merge pull request #4269 from matrix/user_options_bench_info
do not allow --benchmark and --backend-info
2025-06-27 22:02:47 +02:00
Jens Steube
974934dcdf Trying out a tweak to autotune behavior related to -u loop tuning.
Since loop values increase by doubling in autotune, a slow hash-mode
with, for example, 1000 iterations can end up with a suboptimal -u count.
Currently, autotuning starts at 1 and doubles (2, 4, 8, ..., 512, 1024).
If the maximum is 1000, autotune stops at 512, resulting in two kernel
calls: one with 512 iterations and another with 488.

The tweak attempts to find the smallest factor that, when repeatedly
doubled, reaches the target exactly.  For 1000, this would be 125
and for 1024, it would be 1.

However, this logic doesn’t align well with how hashcat handles slow
hash iterations. For instance, PBKDF2-based plugins typically set the
iteration count to N-1, since the first iteration is handled by the
`_init` kernel. So, a plugin might set 1023 instead of 1024, and in such
cases, the logic would incorrectly assume 1023 is the minimum factor
which leads to suboptimal tuning.

To work around this, the factor-finder is executed twice: once with
the original iteration count and once with `iteration count + 1`.
The configuration that results in a lower starting point is used.

Other stuff:

- Fixed a critical bug in the autotuner

This bug was introduced a few days ago. The autotuner has the ability
to overtune the maximum allowed thread count under certain conditions.
For example, in unoptimized -a 0 cracking mode when using rules.
Several parts of the hashcat core require strict adherence to this limit,
especially when shared memory is involved.
To resolve this while retaining overtuning for compatible modes,
a new attribute `device_param->overtune_unfriendly` was introduced.
When set to true, it prevents the autotuner from modifying
`kernel_threads_max` and `kernel_accel_max`.
Four sections in `backend.c` have been updated to set this flag,
though additional areas may also require it.

- Moved the code that aligns `kernel_accel` to a multiple of the compute
  unit count into the overtune section.

- Fixed a bug in the HIP dynloader. It now reports actual error strings,
  provided the API returns them.
2025-06-27 21:52:57 +02:00
Gabriele Gristina
fcc2844880
do not allow --benchmark and --backend-info 2025-06-26 22:12:50 +02:00
Gabriele Gristina
0869e7c1bb
change KERNEL_FQ to KERNEL_FQ KERNEL_FA statements 2025-06-26 21:55:07 +02:00
Jens Steube
bdc47abbe0
Merge pull request #4268 from matrix/modules_unstable_warning_fix
restore module_unstable_warning, excluding AMD Radeon
2025-06-26 21:44:31 +02:00
Gabriele Gristina
a3afca56b8
Merge remote-tracking branch 'upstream/master' into ripemd320 2025-06-26 21:37:40 +02:00
Gabriele Gristina
a5fd75f6ab
restore module_unstable_warning, excluding AMD Radeon 2025-06-26 21:35:11 +02:00
Jens Steube
3182af1bc9 - Renamed shuffle() in inc_hash_scrypt.cl to avoid name collision with
shuffle() present in some OpenCL runtimes
- Updated autotune logic: if the best kernel-loop is not yet found and
  the current kernel-loops setting resulting in a kernel runtime which
  is already above a certain threshold, do not skip to kernel-threads
  or kernel-accel section if no variance is possible
- Revised all plugin module_unstable_warning() checks for
  AMD Radeon Pro W5700X GPU on Metal: rechecked with the latest
  Metal version and removed those now fixed
- Inform the user on startup when backend runtimes and devices are
  initialized
- Fixed some file permissions in the tools/ folder
2025-06-26 19:36:06 +02:00
Jens Steube
825491aa6c Rewrote the SCRYPT core to work around a segmentation fault bug in the Intel OpenCL CPU runtime, without negatively affecting other runtimes.
Fixed the automatic kernel acceleration adjustment routine to account for some OpenCL runtime's buffer size limitation (1/4).
Added a missing license header to scrypt_commit.c (MIT).
2025-06-26 09:47:36 +02:00
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
62a5a85dd6 Added 'next_power_of_two()' and moved both 'next_power_of_two()' and 'previous_power_of_two()' to 'shared.c'
Improved autotuner tweak logic and added boundary checks for accel and threads
Fixed available host memory detection on Windows
Fixed compilation error in MSYS2 native shell
Introduced an 8 GiB host memory usage limit per GPU, even if more is available
Replaced fixed-size host memory detection per GPU with a dynamic kernel-accel based method (similar to GPU memory detection)
Disabled hash-mode autodetection in the python bridge
Removed default invocation of 'rocm-smi' in 'benchmark_deep.pl' to avoid skewed initial results
Reduced default runtime in 'benchmark_deep.pl' scripts due to improved benchmark accuracy in hashcat in general
2025-06-25 11:21:51 +02:00
Jens Steube
669bd619cd
Merge pull request #4267 from lhywk/fix/socket-leak-brain-server
Fix socket leak in brain_server()
2025-06-25 07:24:45 +02:00
lhywk
45a3820560 Fix: close server_fd on brain_server() 2025-06-25 04:47:03 +00:00
Jens Steube
189e8acdf9 Small tweak for new autotuner 2025-06-24 22:21:47 +02:00
Jens Steube
69a585fa4a Autotune refactoring II: dynamic threads-per-block
- Integrated occupancy hints from vendor APIs (CUDA, HIP) to set a
  dynamic threads-per-block limit per kernel instead of using static
  values.
- Added `find_tuning_function()` to identify the relevant kernel.
- Autotuner now runs in three stages: threads -> loops -> accel. The
  first two stages now stop increasing when the tested kernel runtime
  gets too close to the target runtime (96ms for `-w 3`), leaving
  headroom for the next stage to adjust in a finer sense.
- Accel tuning now uses a capped floating-point multiplier instead of
  powers of two.
- Removed workarounds for missing thread autotuning in plugins.
- Removed the hardcoded 4GiB host memory limit for accel. Added a
  cross-platform `get_free_memory()` to check actual free RAM during GPU
  initialization, preventing underutilization of high-end GPUs like the
  4090. If needed, users can still cap memory usage with `-T` or `-n`.
- Updated enums for ROCm 6.4.x and CUDA 12.9.
- Added code to detect kernel register spilling. That's relevant so we
  can keep free enough global memory on the runtime for the runtime to
  handle spills efficiently.
2025-06-24 20:19:42 +02:00
Jens Steube
b241439cf0
Merge pull request #4251 from matrix/module_unstable_warning_AMD-EPYC
Modules: Updated module_unstable_warning for hash-mode 1500, 3000, 14000, 19200
2025-06-23 22:41:54 +02:00
Jens Steube
fd98bbb075
Merge pull request #4263 from matrix/memory_leaks_2025
fix memory leaks
2025-06-23 22:32:08 +02:00
Jens Steube
13a7b56feb Improve the logic for when to use funnelshift and when not to. Some algorithms, such as SHA1-HMAC and DCC1, do not work well with it, so it has been disabled for them.
Fix the automatic reduction of the kernel-accel maximum based on available memory per device by accounting for the additional size needed to handle register spilling.
Fix the tools/benchmark_deep.pl script to recognize benchmark masks more reliably.
2025-06-23 12:30:12 +02:00
Gabriele Gristina
baea7933a8
fix memory leaks (2) 2025-06-22 22:04:38 +02:00
Gabriele Gristina
0720d20cf3
fix memory leaks 2025-06-22 21:47:28 +02:00
Jens Steube
957f3b8ac9
Merge pull request #4262 from matrix/is_autotune_metal
fix is_autotune checks on Apple Metal
2025-06-22 21:14:02 +02:00
Gabriele Gristina
df30dfd5c4
fix is_autotune checks on Apple Metal 2025-06-22 21:06:31 +02:00
Jens Steube
13f48e563d
Merge pull request #4109 from matrix/makefile-clangIdentification
Fixed clang identification in src/Makefile
2025-06-22 20:29:50 +02:00
Jens Steube
ed10e6a913 Autotune and Benchmark refactoring
This change affects three key areas, each improving autotuning:

- Autotune refactoring itself

The main autotune algorithm had become too complex to maintain and has
now been rewritten from scratch. The engine is now closer to the old
v6.0.0 version, using a much more straightforward approach.

Additionally, the backend is now informed when the autotune engine runs
its operations and runs an extra invisible kernel invocation. This
significantly improves runtime accuracy because the same caching
mechanisms which kick in normal cracking sessions now also apply during
autotuning. This leads to more consistent and reliable automatic
workload tuning.

- Benchmarking and '--speed-only' accuracy bugs fixed

Benchmark runtimes had become too short, especially since the default
benchmark mask changed from '?b?b?b?b?b?b?b' to '?a?a?a?a?a?a?a?a'. For
very fast hashes like NTLM, benchmarks often stopped immediately when
base words needed to be regenerated, producing highly inaccurate
results.

This issue also misled users tuning '-n' values, as manually
oversubscribing kernels could mask the problem, creating the impression
that increasing '-n' had a larger impact on performance than it truly
does. While '-n' still has an effect, it’s not as significant. With this
fix, users achieve the same speed without needing to tune '-n' manually.

The bug was fixed by enforcing a minimum benchmark runtime of 4 seconds,
regardless of kernel runtime or kernel type. This ensures more stable
and realistic benchmark results, but typically increasing the benchmark
duration by up to 4 seconds.

- Kernel-Threads set to 32 and plugin configuration cleanup

Some plugin configurations existed solely to work around the old
benchmarking bug and can now be removed. For example,
'OPTS_TYPE_MAXIMUM_THREADS' is no longer required and has been removed
from all plugins, although the parameter itself remains to avoid
breaking custom plugins.

Because increasing threads beyond 32 no longer offers meaningful
performance gains, the default is now capped at 32 (unless overridden
with '-T'). This simplifies GPU memory management. Currently, work-item
counts are indirectly limited by buffer sizes (e.g., 'pws_buf[]'), which
must not exceed 4 GiB (a hard-coded limit). This buffer size depends on
the product of 'kernel-accel', 'kernel-threads', and the device’s
compute units. By reducing the default threads from 1024 to 32, there is
now more space available for base words.
2025-06-22 20:17:52 +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
b7c8fcf27c Removed shared-memory based optimization for SCRYPT on HIP, because the shared-memory buffer is incompatible with TMTO, which is limiting SCRYPT-R to a maximum of 8. This change also simplifies the code, allowing removal of large sections of duplicated code. Removed the section in scrypt_module_extra_tuningdb_block() that increased TMTO when there was insufficient shared memory, as this is no longer applicable.
Refactored inc_hash_scrypt.cl almost completely and improved macro names in inc_hash_scrypt.h. Adapted all existing SCRYPT-based plugins to the new standard. If you have custom SCRYPT based plugins use hash-mode 8900 as reference.
Fixed some compiler warnings in inc_platform.cl.
Cleaned up code paths in inc_vendor.h for finding values for HC_ATTR_SEQ and DECLSPEC.
Removed option --device-as-default-execution-space from nvrtc for hiprtc compatibility. As a result, added __device__ back to DECLSPEC.
Removed option --restrict from nvrtc compile options since we actually alias some buffers.
Added --gpu-max-threads-per-block to hiprtc options.
Added -D MAX_THREADS_PER_BLOCK to OpenCL options (currently unused).
Removed all OPTS_TYPE_MP_MULTI_DISABLE entries for SNMPv3-based plugins.
These plugins consume large amounts of memory and for this reason,limited kernel_accel max to 256. This may still be high, but hashcat will automatically tune down kernel_accel if insufficient memory is detected.
Removed command `rocm-smi --resetprofile --resetclocks --resetfans` from benchmark_deep.pl, since some AMD GPUs become artificially slow for a while after running these commands.
Replaced load_source() with file_to_buffer() from shared.c, which does the exact same operations.
Moved suppress_stderr() and restore_stderr() to shared.c and reused them in both Python bridges and opencl_test_instruction(), where the same type of code existed.
2025-06-21 07:09:20 +02:00
Jens Steube
f399c97db0
Merge pull request #4252 from visitorckw/fix-sort-pot-orig-line
Fix incorrect comparison result in sort_pot_orig_line()
2025-06-20 14:56:06 +02:00
Jens Steube
22e2384314
Merge pull request #4253 from visitorckw/fix-sort-by-src-len
Fix incorrect comparison result in sort_by_src_len()
2025-06-20 14:55:54 +02:00
Jens Steube
1f3247be77
Merge pull request #4256 from DhruvTheDev1/patch-1
Fix Typo: change 'generatic' to 'generic' in both hashcat-assimilation-bridge-development.md and usage.c
2025-06-20 14:54:26 +02:00
Dhruv
301dbe9bf3
Update usage.c
fix spaces
2025-06-20 13:16:43 +01:00
Jens Steube
1e9f71990d
Merge pull request #4257 from DhruvTheDev1/patch-2
Update hashcat-assimilation-bridge.md
2025-06-19 23:08:50 +02:00
Jens Steube
59ec58e3a2
Merge pull request #4258 from DhruvTheDev1/patch-3
Update hashcat-python-plugin-development-guide.md
2025-06-19 23:08:40 +02:00
Dhruv
ce17d5271c
Update usage.c
Fix typo: change 'generatic' to 'generic' in usage.c
2025-06-19 12:29:57 +01:00
Jens Steube
4ed91e430f Refactor Python bridge to load source files as scripts instead of modules.
Previously, the bridge loaded user source files as Python modules, which required stripping extensions and could lead to confusing name conflicts.
Now, it loads them as regular Python scripts, which is more intuitive and aligns better with user expectations.
2025-06-18 20:16:08 +02:00
Jens Steube
c033873e4b Update hipDeviceAttribute_t for ROCm 6.x
Add hipDeviceProp_t and bindings for hipGetDeviceProperties(), hipGetDeviceProperties is required to retrieve gcnArchName[].
Add gcnArchName[] to select the correct --gpu-architecture value for a specific device when using hiprtc.
Include sm_major and sm_minor for CUDA and gcnArchName for HIP in the kernel filename hash.
Update nvrtc_options[] and hiprtc_options[] to avoid unused variables, eliminating the use of --restrict as a placeholder and preventing nvrtc from aborting.
Add check_file_suffix() and remove_file_suffix() helper functions.
2025-06-18 18:29:47 +02:00
Dhruv
a472c572f6
Update hashcat-python-plugin-development-guide.md
- Changed some mistakes e.g. duplicates "the the" and improved some clarity too.
2025-06-18 15:23:43 +01:00
Dhruv
1a20e78131
Update hashcat-assimilation-bridge.md
- Fixed typos and improved clarity
2025-06-18 15:06:13 +01:00
Dhruv
40a2dbe5ec
Update hashcat-assimilation-bridge-development.md
I believe the generetic is a typo and should be generic in the help/documentation strings for bridge parameters in the bridge_init() function.
2025-06-18 15:01:27 +01:00
Jens Steube
13245b5563 Add HC_ATTR_SEQ macro to CUDA kernels. It is left empty so that users can optionally add __launch_bounds__ or other launch attributes if needed.
Add MAX_THREADS_PER_BLOCK macro to CUDA kernels. It defaults to 1024 or a lower value if limited by the plugin module or specified via the -T command line option.
For CUDA, lower the C++ dialect from C++17 to C++14 to reduce JIT compile time. Also add support for --split-compile and --minimal flags to further improve NVRTC compile performance.
Remove power-hungry NVIDIA settings and fix missing sudo calls in tools/benchmark_deep.pl.
Remove NEW_SIMD_CODE macro from kernels that do not support SIMD (no u32x).
2025-06-18 10:08:56 +02:00
Jens Steube
e9137e8405
Merge pull request #4255 from matrix/scrypt_metal
porting the new scrypt engine to Apple Metal
2025-06-17 22:26:31 +02:00
Gabriele Gristina
564c560dcb
porting the new scrypt engine to Apple Metal 2025-06-17 22:17:59 +02:00
Jens Steube
7fe091f4a3 Always use the low-level API to query used memory on a device and silently ignore if unavailable.
The workaround using EXTRA_SIZE should no longer be needed, so we disable it for now and monitor for any issues with memory allocation.
Fixed a bug where a scrypt-based algorithm on an API would require only a single work item, resulting in size_extra_buffer4 being zero, which fails on OpenCL since it does not allow zero-byte allocations.
Ignore TMTO increase on low scrypt configurations if the R value is higher than 1.
2025-06-17 19:09:53 +02:00