1
0
mirror of https://github.com/hashcat/hashcat.git synced 2025-07-10 00:28:11 +00:00
Commit Graph

20 Commits

Author SHA1 Message Date
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
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
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
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
ceb5ff5641 The Assimilation Bridge (Framework) 2025-05-29 15:38:13 +02:00
Gabriele Gristina
003579d21b Modules: Updated module_unstable_warning 2025-05-11 15:40:52 +02:00
Rosen Penev
795674c939 fix some const-qual warnings
Signed-off-by: Rosen Penev <rosenp@gmail.com>
2023-08-20 20:55:54 -07:00
Gabriele Gristina
5f7aa3f33d
Merge branch 'master' into fix_26900_module_hash_encode 2023-04-28 19:56:18 +02:00
Gabriele Gristina
756e2e07b1 Fixed bug in 26900 module_hash_encode 2023-04-22 18:48:46 +02:00
Gabriele Gristina
2adf735e51 Fixed bug in input_tokenizer when TOKEN_ATTR_FIXED_LENGTH is used and refactor modules 2023-04-11 19:34:01 +02:00
Jens Steube
1f4d8dd0fd Remove OPTS_TYPE_MAXIMUM_THREADS flag where no longer required by AMD HIP jit compiler 2022-11-14 17:41:22 +01:00
philsmd
d9749e8799
change interface, add module_benchmark_charset () 2022-07-15 17:17:57 +02:00
Jens Steube
6fce6fb3ff Update all existing modules to use the stock module marker 2022-04-08 14:11:50 +02:00
Jens Steube
5015bc0d2e Module Parser: Renamed struct token_t to hc_token_t to avoid naming conflict with token_t on MacOS 2021-12-20 13:19:40 +01:00
Jens Steube
5b4ac09e91 User Options: Add new module function module_hash_decode_postprocess() to override hash specific configurations from command line 2021-11-28 13:58:27 +01:00
Jens Steube
93ba57f183 Update more module with OPTS_TYPE_MAXIMUM_THREADS 2021-11-14 10:11:53 +01:00
Gabriele Gristina
c538d5d7ba fix OPTS_TYPE in SNMPv3 HMAC-SHA384-256 / HMAC-SHA512-384 modules 2021-08-12 23:37:07 +02:00
Jens Steube
01738fafa0 Deprecated Plugins: Add new module function module_deprecated_notice() to mark a plugin as deprecated and to return a free text user notice
Added option --deprecated-check-disable to enable deprecated plugins
2021-08-10 17:59:52 +02:00
Gabriele Gristina
87f7bc59e3 reduced the amount of memory allocated in gpu for SNMPv3 HMAC-SHA384-256 2021-08-03 21:29:07 +02:00
Gabriele Gristina
357c23d7a1 Added hash-mode: SNMPv3 HMAC-SHA384-256 2021-07-27 05:17:26 +02:00