Reduced the fixed memory reservation size from 1GiB to 64MiB as a result.
Added a warning when the user sets a thread count on the command line higher than recommended by the runtime (based on available registers and shared memory).
Added host-side logic to detect true funnel shift support and disable kernels using it if not supported on the device.
Updated more plugins to limit register count to 128 on NVIDIA GPUs.
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.
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
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).
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>'.
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.
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.
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.
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).
Do not adjust kernel-accel or scrypt-tmto for GPUs with unified memory, typically integrated GPUs in CPUs (APUs).
Redesign the "4-buffer" strategy to avoid overallocation from naive division by four, which can significantly increase memory usage for high scrypt configurations (e.g., 256k:8:1).
Update the scrypt B[] access pattern to match the new "4-buffer" design.
Allow user-specified kernel-accel and scrypt-tmto values, individually or both, via command line and tuning database. Any unspecified parameters are adjusted automatically.
Permit user-defined combinations of scrypt-tmto and kernel-accel even if they may exceed available memory.
Add support for lower iteration counts per kernel invocation than the default, enabling TMTO for low scrypt configurations, such as N=1024
Use TMTO 2 if it reaches 4 times the device processor count, instead of TMTO 1 always
Improve performance for low scrypt configurations (hash-mode 9300)
Fix unit test for 15700 with correct scrypt configurations
Disable CPU over subscription for SCRYPT based algorithms
Remove existing tuningdb entries due to salsa_r() core
refactor. Update tuningdb engine to prefer file entries,
when available, over automatic discovery.
Improve memory-free detection per device, default
--backend-device-keepfree is now set to 0.
Old brute-force OpenCL behavior can be restored using
--backend-device-keepfree 100.
- Replace V_ALIGNBIT_B32 with funnelshift on HIP
- Improve RC4 performance by preventing inlineing
- Fix leftover code in yescrypt-platform.c
- Update docs/hashcat-assimilation-bridge-development.md
- Only initialize hwmon on host for virtualized backends
- Improve SCRYPT tunings on AMD RX6900XT
- Skip memory-free detection on MS OpenCL platform to avoid crashes
- Improve salt usage of 70100/70200, use decoder/kernels from 8900
- Add REPLACE bridge type support (eg. BRIDGE_TYPE_REPLACE_LOOP)
- Switch 70000, 70100 and 70200 to BRIDGE_TYPE_REPLACE_LOOP
- Add synchronization barriers on d2h copy when using bridges
- Improve speed status display updates when using bridges
- Set AMD_DIRECT_DISPATCH=0 to reduce CPU burning loop on AMD backends
- Set benchmark/selftest hash on 70100/70200 to 16:8:1