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>'.
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
- 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.
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.
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.
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.
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.
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 hard-coded SCRYPT N, R, and P values in modules, except where they are intentionally hardcoded.
Fix a bug that always caused a TMTO value of 1, even when it was not needed.
Respect device_available_mem and device_maxmem_alloc values even if a reliable low-level free memory API is present, and always select the lowest of all available limits.
Fix benchmark_deep.pl mask to avoid UTF-8 rejects.
Improve error messages when the check verifying that all SCRYPT configuration settings across all hashes are identical is triggered.
Also improve the error message shown when the SCRYPT configuration of the self-test hash does not match that of the target hash.
Fix a bug where a low-tuned SCRYPT hash combined with a TMTO could result in fewer than 1024 iterations, which breaks the hard-coded minimum of 1024 iterations in the SCRYPT kernel.
General:
The logic for calculating the SCRYPT workload has been moved
from module_extra_buffer_size() to module_extra_tuningdb_block().
Previously, this function just returned values from a static
tuning file. Now, it actually computes tuning values on the fly
based on the device's resources and SCRYPT parameters. This
was always possible, it just wasn't used that way until now.
After running the calculation, the calculated kernel_accel value
is injected into the tuning database as if it had come from a
file. The tmto value is stored internally.
Users can still override kernel-threads, kernel-accel, and
scrypt-tmto via the command line or via tuningdb file.
module_extra_tuningdb_block():
This is now where kernel_accel and tmto are automatically
calculated.
The logic for accel and tmto is now separated and more
flexible. Whether the user is using defaults, tuningdb entries, or
manual command line overrides, the code logic will try to make
smart choices based on what's actually available on the device.
First, it tries to find a kernel_accel value that fits into
available memory. It starts with a base value and simulates
tmto=1 or 2 (which is typical good on GPU).
It also leaves room for other buffers (like pws[], tmps[], etc.).
If the result is close to the actual processor count,
it gets clamped.
This value is then added to the tuning database, so hashcat can pick
it up during startup.
Once that's set, it derives tmto using available memory, thread
count, and the actual SCRYPT parameters.
module_extra_buffer_size():
This function now just returns the size of the SCRYPT B[] buffer,
based on the tmto that was already calculated.
kernel_threads:
Defaults are now set to 32 threads in most cases. On AMD GPUs,
64 threads might give a slight performance bump, but 32 is more
consistent and reliable.
For very memory-heavy algorithms (like Ethereum Wallet), it
scales down the thread count.
Here's a rough reference for other SCRYPT-based modes:
- 64 MiB: 16 threads
- 256 MiB: 4 threads
Tuning files:
All built-in tuningdb entries have been removed, because they
shouldn’t be needed anymore. But you can still add custom entries
if needed. There’s even a commented-out example in the tuningdb
file for mode 22700.
Free memory handling:
Getting the actual amount of free GPU memory is critical for
this to work right. Unfortunately, none of the common GPGPU APIs
give reliable numbers. We now query low-level interfaces like
SYSFS (AMD) and NVML (NVIDIA). Support for those APIs is in
place already, except for ADL, which still needs to be added.
Because of this, hwmon support (which handles those low-level
queries) can no longer be disabled.
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.
- 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
