arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-21 23:58:07 +00:00
Code Issues Releases Wiki Activity
9,563 Commits 1 Branch 27 Tags 191 MiB
master
Commit Graph

1519 Commits

Author SHA1 Message Date
Jens Steube
53f6693495 Temporary enable HIP 4.4/ROCM 4.5 on Linux and globally set native thread count 2021-11-10 19:32:54 +01:00
Jens Steube
ae8e52f127 Device Threads: Set back default maximum device thread number to 1024 2021-11-06 20:18:01 +01:00
Jens Steube
d3ae131128 Device Threads: The default maximum device thread number has been reduced from 1024 to 64, users can still overwrite with the -T option 2021-11-04 19:29:34 +01:00
Jens Steube
b768f96228 Device Threads: The default maximum device thread number has been reduced from 1024 to 256, users can still overwrite with the -T option 2021-11-03 09:08:17 +01:00
Jukka Ojanen
3cacbe0e1f Add support for xz compressed files 2021-08-25 16:56:55 +03:00
Jukka Ojanen
3bc7b6af90 Avoid directly accessing file handle 2021-08-21 14:29:10 +03:00
Jens Steube
50e28ff306
Merge pull request #2926 from jtojanen/master 
Code cleanup and small fixes
 2021-08-19 14:59:16 +02:00
Jukka Ojanen
fed32a58c7 Always cleanup and initialize all parts of HC_FILE structure 2021-08-18 19:28:00 +03:00
Jukka Ojanen
f3cd32bffd Avoid memory copy in rebuild_pws_compressed_append() and assign instead 2021-08-18 17:46:39 +03:00
Jens Steube
76e388ea47 Memory Management: Check in OpenCL that CL_DEVICE_HOST_UNIFIED_MEMORY is true and if so, then half of the available memory 
NEO driver: Activate the NEO driver (Intel GPU) after it passed several tests in hashcat using the latest NEO driver version
 2021-08-18 16:10:51 +02:00
Jens Steube
9254603960 Fixed out-of-boundary read in PKZIP masterkey kernel if the password candidate has length zero. 
Now that kernel threads are no longer fixed over the entire session, hardware_power and hardware_power_all
needs to be updated the same way as kernel_power and kernel_power_all for each call to inner2_loop().
 2021-08-17 17:38:45 +02:00
Jens Steube
05f107cbe5 Added option --generate-rules-func-sel to specify a pool of operators that the random rule generator is allowed to pick from 2021-08-14 11:43:32 +02:00
Jens Steube
6b485dd3fd
Merge pull request #2934 from jtojanen/fix_extra_tmp_size 
Assume hashconfig->tmp_size is the element size
 2021-08-14 10:06:55 +02:00
Jukka Ojanen
5c0820b8d3 Assume hashconfig->tmp_size is the element size 2021-08-13 13:13:12 +03:00
Jens Steube
1f22984313 The hc_bytealign_be_S() was merged too early, but is not yet used anywhere for this platform so we can remove it again. 2021-08-12 15:15:32 +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
Jens Steube
63a0c5a1e1
Merge branch 'master' into snmpv3_sha512 2021-08-04 21:05:42 +02:00
Jens Steube
dde75f2763
Merge pull request #2904 from matrix/snmpv3_sha384 
Added hash-mode: SNMPv3 HMAC-SHA384-256
 2021-08-04 21:04:55 +02:00
Jens Steube
cb69e2d413 Added some HIP version checks, fall back to OpenCL automatically 
Switched HIP version check from driverVersion to runtimeVersion
Fixed syntax check of HAS_VPERM macro in several kernel includes causing invalid error message for AMD GPUs on Windows
Updated AMD driver requirements
Updated docs/changes.txt with missing changes from previous commits
Fixed invalid vector data type in Murmur Hash in -a 3 mode
Fixed uninitialized variable warning in src/hashes.c
Fixed broken support for --generate-rules-func-min
 2021-08-04 20:49:22 +02:00
Jens Steube
af5d346244
Merge pull request #2894 from jtojanen/master 
Allow async execution
 2021-08-02 10:05:20 +02:00
Jens Steube
20a7b9f992 Tuning-Database: Add new module function module_extra_tuningdb_block() to extend hashcat.hctune content from a plugin 
See src/modules/module_08900.c as an example
 2021-08-01 16:25:37 +02:00
Jukka Ojanen
3d7ce7162b Merge branch 'master' of https://github.com/hashcat/hashcat 2021-08-01 15:36:52 +03:00
Jens Steube
f4e52ca533 Add new rule function '3' to switch the case of the first letter after occurrence N of char X 2021-08-01 00:04:10 +02:00
Jukka Ojanen
c3195d0603 Merge branch 'master' of https://github.com/hashcat/hashcat 2021-07-30 11:34:25 +03:00
Jukka Ojanen
d7cc8d7cd2 Revert due to module_03200.c, module_25600.c and module_25800.c using device_param->kernel_dynamic_local_mem_size_memset 2021-07-28 19:49:02 +03:00
Jukka Ojanen
8b590f651b Remove unused gpu_memset and its references 2021-07-28 19:26:30 +03:00
Jens Steube
ae39fa0979
Merge pull request #2908 from hashcat/master 
Backports
 2021-07-28 07:55:11 +02:00
Jukka Ojanen
5dbe4958b8 Remove digests_shown_tmp buffer, used in check_cracked() to zero out hashes 2021-07-27 19:01:13 +03:00
Jukka Ojanen
cdf27a1cb3 Implement async run_cuda_kernel_memset() and run_cuda_kernel_memset32() 2021-07-27 18:56:59 +03:00
Jukka Ojanen
be33ccaa6c check_hash() needs to return value to indicate possible error 2021-07-27 14:36:48 +03:00
Jens Steube
532a154542 ADL: Updated support for AMD Display Library to 15.0, updated datatypes and added support for OverDrive 7 and 8 based GPUs 2021-07-27 12:02:27 +02:00
Jukka Ojanen
d7de3550b1 Add fast event for CUDA and HIP, and use it to synchronize check_hash() 2021-07-27 12:57:26 +03:00
Gabriele Gristina
357c23d7a1 Added hash-mode: SNMPv3 HMAC-SHA384-256 2021-07-27 05:17:26 +02:00
Gabriele Gristina
15b90c953e Added hash-mode: SNMPv3 HMAC-SHA512-384 2021-07-27 05:14:09 +02:00
Jukka Ojanen
a642f7b233 Remove synchronous GPU memory copy functions 2021-07-26 15:36:42 +03:00
Jens Steube
b66c093c17
Merge pull request #2902 from hashcat/master 
Backports
 2021-07-26 08:18:35 +02:00
Jens Steube
b53691c8f5 ADL: Updated support for AMD Display Library to 14.0, updated datatypes and added support for OverDrive 7 and 8 based GPUs 2021-07-26 07:48:56 +02:00
Jukka Ojanen
4bd7363674 Add missing HIP declarations 2021-07-25 22:18:16 +03:00
Jens Steube
72e307fbce
Merge pull request #2901 from hashcat/master 
Backport changes
 2021-07-25 10:57:39 +02:00
Jens Steube
959a232828
Merge pull request #2885 from neheb/charfixes 
const and char fixes
 2021-07-25 10:36:40 +02:00
Jens Steube
84a4058edf
Merge pull request #2900 from hashcat/master 
Backport changes
 2021-07-25 10:34:05 +02:00
Jens Steube
640d95a00f Vendor Detection: Add "Intel" as a valid vendor name for GPU on macOS 2021-07-24 16:01:30 +02:00
Jens Steube
5ffcaa980d HIP Backend: Added support to support HIP 4.4 and later, but added check to rule out older versions because they are incompatible 2021-07-23 16:04:34 +02:00
Rosen Penev
14f5a26027 use const char for fopen mode 
Fixes -Wwrite-strings warnings.

Signed-off-by: Rosen Penev <rosenp@gmail.com>
 2021-07-22 18:11:11 -07:00
Jens Steube
f9e74045b5 Commandline: Throw an error if separator character given by the user with -p option is not exactly 1 byte 2021-07-22 18:54:02 +02:00
Jukka Ojanen
fafd24237d Define HC_ALIGN macro to control data alignment and use common constant zero buffer in run_cuda_kernel_bzero(), run_hip_kernel_bzero() and run_opencl_kernel_bzero(). 2021-07-22 14:24:03 +03:00
Jens Steube
5024865d87 Kernel Threads: Use warp size / wavefront size query instead of hardcoded values as base for kernel threads 
Kernel Cache: Add kernel threads into hash computation which is later used in the kernel cache filename
Remove some unused function symbol lookups in HIP library
 2021-07-22 11:46:47 +02:00
Jens Steube
c990e252d3 Added option --multiply-accel-disable (short: -M) to disable multiply the kernel-accel with the multiprocessor count automatism 2021-07-21 15:47:05 +02:00
Jukka Ojanen
8674e23d79 Add async HIP memcpy functions: hc_hipMemcpyDtoDAsync(), hc_hipMemcpyDtoHAsync() and hc_hipMemcpyHtoDAsync(). Implement partially async HIP memset and bzero kernels. 2021-07-20 12:47:10 +03:00
Jukka Ojanen
4263cafdcf Add async CUDA memcpy functions: hc_cuMemcpyDtoDAsync(), hc_cuMemcpyDtoHAsync() and hc_cuMemcpyHtoDAsync(). Implement partially async CUDA memset and bzero kernels. 2021-07-20 12:23:39 +03:00
Jukka Ojanen
ea5180ac46 Include missing bzero kernel parameters 2021-07-20 11:59:44 +03:00
Jukka Ojanen
a2a1d04bcf Implement gpu_bzero 2021-07-17 19:00:10 +03:00
Jens Steube
1ab56afb9e
Merge pull request #2874 from hashcat/master 
Backports from master
 2021-07-12 08:57:50 +02:00
Jens Steube
ca3beacd93 Disable dynamic shared memory on HIP, because hipFuncSetAttribute() maps to cudaFuncSetAttribute() and not to cuFuncSetAttribute() 2021-07-11 14:30:49 +02:00
Jens Steube
92cf955ae6
Merge pull request #2872 from jtojanen/master 
Use critical section instead of mutex in Windows
 2021-07-11 13:43:11 +02:00
Jens Steube
1b84a9e53b Add missing backports from code base v6.2.2 
Fix context to thread management
Fix missing code in selftest.c, autotune.c, hashes.c, dispatch.c and backend.c
Use IS_HIP depending code makes it easier for future optimization related to inline assembly calls - instead of using IS_CUDA || IS_HIP
See TODO markers for more optimizations / next steps
 2021-07-11 12:38:59 +02:00
Jens Steube
a22f8149fc
Merge branch 'HIP' into hip 2021-07-10 21:34:09 +02:00
Jukka Ojanen
2fd0a81f93 Use critical section instead of mutex in Windows 2021-07-10 16:29:40 +03:00
Jens Steube
9fc6c26f8c Added temperature watchdog and utilization for CPU on linux using sysfs and procfs 2021-07-10 14:24:51 +02:00
Jens Steube
bfe83ec138 Added temperature watchdog for CPU on linux using sysfs 2021-07-10 08:43:15 +02:00
Jens Steube
979f9e9868 Rename hardware monitor interface sysfs to sysfs_amdgpu 2021-07-09 20:48:10 +02:00
reger-men
ea7b74389f First draft HIP Version 2021-07-09 03:50:40 +00:00
Jens Steube
f146a05878 Added option --markov-inverse to inverse markov statistics, with the idea of reversing the order of the password candidates 2021-07-07 10:32:08 +02:00
Jens Steube
701ad7c441 Fix segfault in --hash-info by avoid huge stack buffer allocation 2021-06-30 21:01:55 +02:00
Jens Steube
56c2243dfb KeePass: Increase supported size for KeePass 1 databases from 300kB to 16MB 2021-06-30 13:16:03 +02:00
Jens Steube
254e33c473 File handling: Do not abort on seeing a BOM in input files, just warn and ignore the BOM 2021-06-29 20:42:22 +02:00
Jens Steube
2af45cd03f Rename hc_fopen_nozip() to hc_fopen_raw() 2021-06-29 20:13:11 +02:00
Jens Steube
dcaba1f473 Brain: Add brain_ctx_t to hashcat_ctx_t to enable runtime check if hashcat was compiled with brain support 2021-06-26 20:24:45 +02:00
Jens Steube
b30d55c5d9 Removed LZMA "Windows.h" capitalized file name workaround used to cross compile Windows binaries from Linux hosts. 
This breaks compiling windows binaries if the underlaying filesystem is not case sensitive. An example would be using /mnt/c folder under WSL.
Any update on the LZMA SDK needs to be post processed with this command: sed -i 's/<Windows.h>/<windows.h>/' deps/LZMA-SDK/C/*.c deps/LZMA-SDK/C/*.h
 2021-06-25 20:30:26 +02:00
Jens Steube
90f0e78b5b Add event about autodetection process start 
Remove "..." substring from final messages
Add patient message on "..." startup messages
Add missing docs/changes.txt entry
 2021-06-22 11:37:59 +02:00
Jens Steube
7e267b9b37
Merge pull request #2825 from matrix/hwmon_osx_v2 
Add OSX HW Monitor initial support (2021)
 2021-06-20 11:09:35 +02:00
Gabriele Gristina
3a6481da4e show all fans in hwmon output 2021-06-19 20:32:07 +02:00
Jens Steube
3119525ea3 Fixed integer overflow in Recovered/Time status view column caused by division > 0 but < 1 2021-06-18 07:57:15 +02:00
Jens Steube
62397283c1 VeraCrypt: Increase password length support for non-boot volumes from 64 to 128. 
See https://github.com/hashcat/hashcat/issues/2616 for details.
 2021-06-17 09:45:30 +02:00
Gabriele Gristina
4dc20533bd fix ext_iokit.h 2021-06-12 20:24:43 +02:00
Gabriele Gristina
bc4ce4cbeb Add support for CPU/GPU device temperature and fanspeed using iokit (Apple) 2021-06-12 20:13:31 +02:00
Jens Steube
2b68d4bec1
Merge pull request #2821 from matrix/autodetect-only 
add identify/autodetect-only option
 2021-06-12 10:26:33 +02:00
Gabriele Gristina
340e8d1033 fix identify functionality 2021-06-11 19:25:39 +02:00
Gabriele Gristina
3c6f4b0e69 add identify/autodetect-only option 2021-06-10 20:44:21 +02:00
Gabriele Gristina
e5de44c80d Restyling hwmon (2021)(2) 2021-06-10 20:13:12 +02:00
Bjoern Kerler
d59f8c42d2 Prepare FBE support 2021-06-10 10:31:27 +02:00
Jens Steube
186e68289c
Merge pull request #2804 from matrix/autodetect 
Added new feature: autodetect hash-mode
 2021-06-06 18:20:44 +02:00
Gabriele Gristina
b8ae1444de skip autodetect if benchmark is enabled, keep table sorting like usage 2021-06-06 13:40:01 +02:00
Jens Steube
c9db099330
Merge pull request #2456 from nycex/xdg 
use XDG_DATA_HOME for profile_dir
 2021-06-06 10:51:53 +02:00
Gabriele Gristina
ced9b0c6f8 Added new feature: autodetect hash-mode 2021-06-05 22:11:18 +02:00
nycex
40c68b8bf4
use recursive mkdir for the profile and the cache dir 2021-06-05 21:55:59 +02:00
Jens Steube
72d1937fe7 Modules: Added hash categories HASH_CATEGORY_IMS and HASH_CATEGORY_CRYPTOCURRENCY_WALLET 2021-06-05 15:32:28 +02:00
nycex
470e844e5d
use XDG_CACHE_HOME for kernels 2021-06-05 00:38:43 +02:00
Jens Steube
f54643479d CRC32: Prevent decompression of data used in CRC32 calculation on host. This leads to false negatives with TrueCrypt/VeraCrypt keyfiles 2021-05-29 13:27:05 +02:00
Jens Steube
0b52409590 Modules: Recategorized HASH_CATEGORY option in various modules 2021-05-28 20:06:01 +02:00
Jens Steube
9f5a22a3ab OpenCL Backend: Use CL_DEVICE_BOARD_NAME_AMD instead of CL_DEVICE_NAME for device name in case OpenCL runtime supports this query 2021-05-26 19:09:50 +02:00
Jens Steube
5c6501444a Kernels: Add standalone true UTF8 to UTF16 converter kernel that runs after amplifier. Use OPTS_TYPE_POST_AMP_UTF16LE from plugin 2021-05-20 14:34:24 +02:00
Jens Steube
307cb63661
Merge pull request #2695 from matrix/cryptoapi 
Added Linux Kernel Crypto API (2.4) KPA
 2021-05-17 23:41:10 +02:00
Jens Steube
e365313a44 Dependencies: Updated LZMA SDK from 19.00 to 21.02 alpha 2021-05-16 08:51:59 +02:00
Jens Steube
bb402b784a Update module_unstable_warning for benchmark short selection on macOS for CPU and GPU; Allow use of GPU without --force testwise 2021-05-10 14:36:41 +02:00
Gabriele Gristina
17a83a50af cryptoapi opts (1) 2021-05-06 23:24:46 +02:00
Jens Steube
32b40d5c53 Added command prompt [f]inish to instruct hashcat to finish the current attack and quit afterwards 2021-05-05 22:53:04 +02:00
Jens Steube
0c2afde83b Add support for clUnloadPlatformCompiler() 2021-05-02 08:15:25 +00:00
Jens Steube
59459d0e5b Fixed memory leak causing problems in sessions with many iterations. for instance, --benchmark-all or large mask files 2021-05-01 17:27:33 +02:00
Konrad Goławski
3535f3ae55 Change typedef hashconfig_t 2021-04-27 19:09:44 +02:00
Jens Steube
15f35fa68c Scrypt Kernels: Reduced kernel wait times by making it a true split kernel where iteration count = N value 2021-04-21 15:59:14 +02:00
Jens Steube
56f47cabe2 Fixed race condition in potfile check during removal of empty hashes 2021-04-21 09:22:00 +02:00
Jens Steube
62fc3601bb Wrap atomic functions with hc_ prefix to have better platform control 2021-04-20 17:47:44 +02:00
Jens Steube
57a8923b81 Update complete SCRYPT workload tuning logic. 
A detailed description will follow.
Set -m 8900 defaults to 16k:8:1 (default scrypt settings).
 2021-04-16 20:17:53 +02:00
Jens Steube
ff96015f53 Add OPTS_TYPE_NATIVE_THREADS for use by plugin developer to enforce native thread count (useful for scrypt) 2021-04-14 15:22:30 +02:00
Jens Steube
9033975efd Allow plugins to disable the multiplication of the kernel-accel value with the multiprocessor count of the compute device. Will be used later. 2021-04-13 11:26:17 +02:00
Jens Steube
a0eae9050c OpenCL Runtime: Workaround JiT compiler deadlock on NVIDIA driver >= 465.89 2021-04-11 13:35:40 +02:00
Jens Steube
b7951262e7
Merge pull request #2590 from philsmd/include_header_fix 
Fix some include headers macros
 2021-04-07 20:47:20 +02:00
Jens Steube
f7c3a907a2
Merge pull request #2585 from ventaquil/types-guard 
Fix types.h include guard
 2021-04-07 20:47:05 +02:00
Jens Steube
1ba1fa6e58
Merge branch 'master' into full-aes-gcm-and-hash-mode-27000 2021-04-05 17:51:43 +02:00
Jens Steube
1323ef3a82 Support loading hash from command line in -m 22000 and -m 22001 
Fixes #2742
 2021-04-04 11:38:02 +02:00
Gabriele Gristina
77e328d659 Removed option --example-hashes, now is an alias of --hash-info 2020-12-29 07:56:20 +01:00
Gabriele Gristina
3ed1f0d840 Added new option: --hash-info 2020-12-29 04:58:58 +01:00
Gabriele Gristina
54df7d53ea Added full AES-GCM cipher & hash-mode 27000 - Stargazer Stellar Wallet XLM 2020-12-25 04:50:18 +01:00
philsmd
a3663d1cc4
Fix some include headers macros 2020-10-23 17:29:39 +02:00
Konrad Goławski
3c251416b0 Fix types.h include guard 2020-10-20 16:27:29 +02:00
Jens Steube
04d5e5a119 New Attack-Mode: Association Attack. Like JtR's single mode. Very early 
stage. See hashcat Forum for detailed writeup.
 2020-09-29 15:56:32 +02:00
Jens Steube
57bef8abc9 Display possible NVIDIA CUDA/RTC library loading error message only in case a NVIDIA device was found using OpenCL 2020-09-06 13:29:32 +02:00
Jens Steube
98aef2ae92 Module Structure: Add 3rd party library hook management functions. This also requires an update to all existing module_init() 2020-08-29 16:12:15 +02:00
Jens Steube
a72ba6faab Add OPTI_TYPE_SLOW_HASH_SIMD_INIT2 and OPTI_TYPE_SLOW_HASH_SIMD_LOOP2 2020-08-14 15:52:36 +02:00
Jens Steube
6cec2664bc Add new category HASH_CATEGORY_PRIVATE_KEY 2020-07-30 12:25:46 +02:00
Jens Steube
61e8368592 hashcat 6.1.1 2020-07-29 12:16:53 +02:00
Jens Steube
9ed07f892b hashcat 6.1.0 2020-07-28 12:29:13 +02:00
Jens Steube
0ff2f8c5e1 OpenCL Devices: Utilize PCI domain to improve alias device detection 2020-07-27 15:21:56 +02:00
Jens Steube
5fab483cf0 Compile ZLIB: Fixed makefile include paths in case USE_SYSTEM_ZLIB is used 2020-07-21 12:41:20 +02:00
Jens Steube
577768320d Backend: Changed the maximum number of compute devices from 64 to 128 2020-07-10 11:32:42 +02:00
Jens Steube
13e2dec017 hashcat 6.0.0 2020-06-16 17:30:11 +02:00
philsmd
437e9aa3d1
allow --left and -o work with binary hash files 2020-06-15 15:43:49 +02:00
Jens Steube
84d3981e7b Add returncode PARSER_HAVE_ERRNO as a valid returncode from module_hash_binary_count() and module_hash_binary_parse() module hooks. You can use this in case fopen() fails and you want to transport back the original error from the OS 2020-05-27 15:23:02 +02:00
Jens Steube
bb19c9cf1b
Merge pull request #2373 from mpalmer/more-error-messages 
Add some more generic error messages
 2020-05-26 13:25:36 +02:00
Jens Steube
ff37264c36 Use special type for hashlist sorting in brain_compute_session() 2020-05-12 08:54:26 +02:00
Jens Steube
f9e4dc0db1 Remove sort_by_string(), no longer needed 2020-05-11 12:00:26 +02:00
Jens Steube
d5cd0a0030 Fix invalid brain session id calculation and move sorting functions to shared.c 2020-05-11 11:52:36 +02:00
Jens Steube
33579ae3db Add new option for modules OPTS_TYPE_SELF_TEST_DISABLE to disable self-test functionality from hash-mode directly 2020-05-08 10:18:17 +02:00
Matt Palmer
b25c8d9a5a Add some more generic error messages 
These are almost all taken from error messages that are specific to one
particular module, and turned into generic errors that can be used for any
module.  It seemed like a better idea to provide generic messages than to
encourage infinite proliferation of error codes (which would probably end up
blowing the 255 value threshold sooner or later).  It doesn't seem necessary
to provide module-specific error messages for things like "Invalid <X>
filesize", since users should already know what sort of file they're asking
to be parsed.
 2020-04-23 12:36:15 +10:00
Jens Steube
ba7163062d Do not set -cl-std=XXX to workaround NEO driver bug causing to hang while compiling -m 22000 2020-03-13 09:43:41 +01:00
Jens Steube
2b2a7ede66 OpenCL Options: Set --spin-damp to 0 (disabled) by default. With the CUDA backend this workaround became deprecated 2020-03-12 10:51:10 +01:00
Jens Steube
d706f90a75 Dictstat: On Windows, the st_ino attribute in the stat struct is not set which can lead to invalid cache hits. Added the filename to the database entry. 2020-03-10 15:42:55 +01:00
Jens Steube
f8f5e1cc93 Remove unused OPTS_TYPE_PREFERED_THREAD 2020-03-06 08:53:43 +01:00
Jens Steube
a4db1a0abd Remove unused OPTS_TYPE_ST_HASH_MD5 2020-03-05 12:43:48 +01:00
Jens Steube
f96e35649d Change bitsliced kernels from 3d to 2d invocation mode for slightly better performance 2020-02-22 07:59:58 +01:00
Jens Steube
d9473358ef Add support for OPTS_TYPE_LOOP_EXTENDED kernel for special cases like VeraCrypt 2020-02-20 16:00:21 +01:00
Jens Steube
9957422191 Add tokenizer option TOKEN_ATTR_SEPARATOR_FARTHEST as an option to deal with hash formats where separator character could be part of the salt data itself and at the same time does not allow escape character logic to be applied. This can only work if it is guaranteed by the design of the hash format that the separator character does not occur after the position of the separator character. 2020-02-19 10:35:44 +01:00
Jens Steube
4ed18af14c Make VeraCrypt PIM configuration mechanism easier to read 2020-02-08 12:24:47 +01:00
Jens Steube
4788c61dd2 Add OPTI_TYPE_REGISTER_LIMIT flag to enable register limiting in CUDA 2020-02-04 21:53:27 +01:00
Jens Steube
c40f474c2e Add special module option to indicate the kernel is using dynamic shared memory 2020-02-02 11:24:38 +01:00
Jens Steube
aef53f7e10 OpenCL Runtime: Allow the kernel to access post-48k shared memory region on CUDA. Requires both module and kernel preparation 2020-02-01 14:27:42 +01:00
Jens Steube
1fc37c25f9 OpenCL Kernels: Moved "gpu_decompress", "gpu_memset" and "gpu_atinit" into new OpenCL/shared.cl in order to reduce compile time 2020-02-01 09:00:48 +01:00
Jens Steube
346637ec43 Improve cujit logging 2020-01-30 11:44:57 +01:00
Jens Steube
66ae5125ce Cache cubin instead of PTX to decrease startup time 2020-01-29 15:56:36 +01:00