only the first hash in a multihash list was marked as cracked, regardless
of which hash was actually cracked. For example, if the second hash was
cracked, it incorrectly marked the first as cracked and left the second
uncracked. This issue only affected beta versions and only in multihash
cracking mode.
Added deep-comp kernel support for Kerberos modes 28800 and 28900,
enabling multihash cracking for the same user in the same domain, even if
the password was changed or the recording was bad.
Added a rule ensuring that device buffer sizes for password candidates,
hooks, and transport (tmps) must be smaller than 1/4 of the maximum
allocatable memory. If not, hashcat now automatically reduces kernel-accel
down to 1, then halves the number of threads and restores kernel-accel up
to its maximum, repeating until the size requirement is met.
Fixed salt length limit verification for -m 20712.
Fixed password length limit for -m 14400.
Fixed unit test salt generator for -m 21100, which could produce duplicate
hashes under certain conditions.
Added the OPTS_TYPE_NATIVE_THREADS flag to the following hash modes
(after benchmarking): 7700, 7701, 9000, 1375x, 1376x, 14800, 19500, 23900.
- Used blake2b_transform() instead of blake2b_update() to avoid compiler problems
on Intel OpenCL and segfaults on POCL (still unsure of exact cause but possibly
related to the shuffle functions in combination with these OpenCL drivers).
- Remove 'bug' comments (these are resolved now).
- Added implementation of 'argon2_hash_block()' for non-warped (CPU) case.
- Introduced 'LBLOCKSIZE' for the size of an argon2 block per thread in u64.
Most of the code should now be able to support any warp/wavefront size.
- Replaced inline asm in hc_byte_perm() with __builtin_amdgcn_perm()
- Replaced inline asm in hc_bytealign() with __builtin_amdgcn_alignbyte()
- Defined HC_INLINE as default for HIP, significantly boosting kernel performance of pure kernels
- Removed IS_ROCM from inc_vendor.h as it's no longer needed
- Removed backend-specific code from several hash-modes and inc_rp_optimized.cl, as hc_bytealign_S() is now available on all backends
This commit introduces initial support for mixed mode multihash cracking
in Argon2. Although I was skeptical at first, the final solution turned
out better than expected with only a minimal speed loss (1711H/s ->
1702H/s).
Unit tests have been updated to generate random combinations of
Argon2-I/D/ID with randomized m, t, and p values. So far, results look
solid.
Note: This is a complex change and may have undiscovered edge cases.
Some optimization opportunities remain. JIT-based optimizations are not
fully removed. We could also detect single-hash scenarios at runtime
and disable self-tests to re-enable JIT. Currently, the kernel workload
is sized based on the largest hash to avoid out-of-bound memory access.
Fixed compiler warnings in inc_hash_argon2.cl.
Moved argon2_tmp_t and argon2_extra_t typedefs from argon2_common.c back to the module to allow plugin developers to modify them when using Argon2 as a primitive.
Slightly improved autotune behavior for edge cases such as 8700 and 18600, where some algorithms started with theoretical excessively high value, leaving no room for proper tuning.
Removed argon2_module_kernel_threads_min() and argon2_module_kernel_threads_max() from argon2_common.c. Switched to using OPTS_TYPE_NATIVE_THREADS instead. Plugin developers can still use it. This simplifies CPU integration, as CPUs typically run with a single thread.
Updated plugins 15500 and 20510. Added a thread limit to prevent autotune from selecting an excessively high thread count. The issue originated from the runtime returning an unrealistically high ideal thread count.
Fixed parameter types in inc_hash_blake2b.cl and inc_hash_blake2s.cl for FINAL value.
Added kernel code for -m 15400 to s04/s08/m04/m08, even if not needed, to help autotune find optimal workitem settings.
Fixed a rare autotune case (e.g. in mode 18600) where threads_min was not a multiple of kernel_preferred_wgs_multiple, and changes it so that as long as it only threads_min is affected and not threads_max, we now ensure at least kernel_preferred_wgs_multiple.
Improved autotune logic for best thread count: double thread count until reaching the device's preferred multiple, then increase in steps of that multiple while comparing efficiency vs. runtime, and select the configuration with best efficiency, not highest thread count.
Always set funnelshift support to true for HIP devices, as it always reports false.
Set minimum loop count to 250 for all VeraCrypt modes with PIM brute-force support.
- added support to 2D/3D Compute
- improved compute workloads calculation
Makefile:
- updated MACOSX_DEPLOYMENT_TARGET to 15.0
Unit tests:
- updated install_modules.sh with Crypt::Argon2
Argon2 start works with Apple Metal
=====================================
This patch modifies the existing Argon2 plugin, which was initially
designed to work only with CUDA. Supporting OpenCL and HIP required
broader architectural changes.
1. The tmps[] structure no longer holds the "large buffer". This
buffer stored the scratch areas for all password candidates in one
chunk. But we do not need to hold scratch areas for all candidates
simultaneously. All we need to do is hold chunks large enough
per password.
To simplify logic, the buffer is not divided by password count, but
divided by four, which fits within the "1/4 global memory" limit on
some OpenCL runtimes.
Hashcat already had logic to support this, but the buffer needed to be
moved to a different buffer type. It has now been relocated from the
"tmp buffer" to the "extra tmp buffer", following the same strategy
used in newer SCRYPT plugins.
This improves handling across several subcomponents:
- Hashcat backend divides into four asymmetric buffers, hence the
name "4-buffer strategy"
- If the candidate count isn't divisible by 4, leftover candidates are
assigned to the first (and possibly second and third) buffer
- No code in the plugin is required, as this was designed for exactly
such cases where future algorithms require a lot of memory
- Plugin was rewritten to report the size needed in
module_extra_tmp_size(), which triggers the "4-buffer" strategy
- The split is not even, but each part is large enough to hold
a multiple of a full scratch buffer for a password
- The kernel code in m34000_init/loop/comp now uses a code block
that finds its buffer by doing "group_id % 4"
- Prevents the need to over-allocate memory to avoid OOB access
- The original "tmps buffer" now holds a small dummy state buffer
2. Replaced warp shuffle instruction
The instruction __shfl_sync() is not available in runtimes
other than CUDA. Some have alternatives, some do not.
To prevent branching per backend runtime, the new general macro
hc__shfl_sync() replaces all calls to __shfl_sync().
This allows us to implement runtime-specific solutions and
take effect at compile time to prevent regressions.
- CUDA:
We simply map to the original __shfl_sync()
- HIP:
We map to shfl(), a built-in intrinsic. This instruction doesn't
support masks like __shfl_sync() does, but masks are not needed
in Argon2 anyway. It requires an additional parameter, the wavefront
size. This is natively 64, but we hardcode this to 32 so it aligns
with NVIDIA's warp size.
- OpenCL:
- AMD: We have access to the instruction __builtin_amdgcn_ds_bpermute().
This instruction only supports 32-bit integers, requiring us to
pack and unpack the 64-bit values manually
- NVIDIA: We use inline assembly with "shfl.sync.idx.b32". Same as
with AMD, we need to pack and unpack 32-bit integers. The 64-bit
support in CUDA is just overloaded and internally does the same thing.
- Others: We use a shared memory pool and combine it with a barrier.
This LOCAL_VK pool must be sized at compile time and transported to
the Argon2 code in "inc_hash_argon2.cl". This required changing all
function declarations that use shuffles slightly.
Unlock full threading for init and comp kernels
===============================================
This is implemented using a new flag:
OPTS_TYPE_THREAD_MULTI_DISABLE
Behavior is similar to:
OPTS_TYPE_MP_MULTI_DISABLE
It simply disables the multiplier normally applied to password batch size.
But attention, this change completely unbinds this effect from the
real threads spawned on the compute device. If the thread count is not
set to 1 in the plugin, it will start autotuning it.
In the case of Argon2, we hard-code it to 32 instead, which also changes
how "warp size" was used in the original implementation, and which is not
compatible with HIP and/or OpenCL. However, we need to maintain this thread
size to utilize warp shuffle and its alternatives in other runtimes.
Benefits:
- Enables full threading for init and comp kernels (1667 H/s to 1722 H/s)
- Allows future algorithms to enable parallel processing of single
password candidates, if supported
Plugin changes:
- Removed the "hack" where thread count = 1 disabled the multiplier
- Removed per-device warp count detection code and struct changes
- Removed warp handling and "num_elements / thread_count" division in
the run_kernel() function
Simplified autotune logic for Argon2
====================================
The goal is to calculate the maximum number of password candidates that
can run in parallel, constrained only by device memory.
- Removed all code related to Argon2 from autotune
- Implemented in "module_extra_tuningdb_block()" (like SCRYPT)
- We create a tuningdb entry at runtime!
- Still allows override via tuningdb or CLI
- Considers register spilling (read at startup)
- Prevents global-to-host memory swap performance issues
Add Argon2I and ArgonD support
==============================
The kernel prepared from NFI already had support for the different Argon
types. No change was needed.
To support the other Argon2 types, the tokenizer had to be improved to
support a variety of different signatures in the same hash-mode.
Bugfixes
========
- Fixed missing entries in "switch_buffer_by_offset_8x4_le_S()"
- Fixed benchmark hash misdetection for scrypt. This was due to
outdated logic used in scrypt to detect whether the plugin was
called from a benchmark session or a regular one
- Fixed a bug in "module_hash_encode()" where Base64 padding '=' was
retained
- Fixed missing "GLOBAL_AS" / "PRIVATE_AS" casts for OpenCL
- Fixed compiler warnings (e.g., "index_u32x4()", "get_group_id()")
by adding return values
- Fixed a bug in token.len_max[6], which was allowing decoding
of a 256-byte data into a 16-byte buffer (digest)
Other improvements
==================
- Added unit test module for automated testing
- Added support to the tokenizer to allow multiple signatures.
Leave out TOKEN_ATTR_FIXED_LENGTH to enable this in your plugins
- Updated "hc_umulhi()", also exists for HIP
- Renamed "gid" to "bid" when using "get_group_id()" for clarity
- Removed "#ifdef IS_CUDA" as all backends are now supported
- Removed deprecated "OPTS_TYPE_MAXIMUM_ACCEL" attribute
Performance note
================
For testing, I used the self-test hash configured according to the
RFC 9106 recommendation: m=65536, t=3, p=1.
In my benchmarks, the AMD RX 7900 XTX achieved 1401 H/s using the same
hash that was used to test NVIDIA's RTX 4090. The RTX 4090 reached
1722 H/s, making it faster in absolute terms. However, at the time of
writing, it is more than three times as expensive as the 7900 XTX.
It's also worth noting that an older NVIDIA GTX 1080 Ti still reached
565 H/s with the same test vector, and may be found at significantly
lower cost.
Across all tested Argon2 configurations, the performance gap between
the RX 7900 XTX and the RTX 4090 remained proportionally consistent,
indicating a clear linear scaling relationship between the two GPUs.
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
