From f7a8e7c54b71e1352df0b6269bee6fa426ad162a Mon Sep 17 00:00:00 2001 From: jsteube Date: Thu, 29 Jun 2017 12:19:05 +0200 Subject: [PATCH] Multiple changes: * Added more preparations to support to crack passwords and salts up to length 256 * Added option --length-limit-disable to disable optimization based on password- and salt-length * Added option --self-test-disable to disable self-test functionality on startup --- OpenCL/inc_scalar.cl | 42 +++++++++ OpenCL/m00000_a1.cl | 98 ++++++++++++++++++++ docs/changes.txt | 3 + extra/tab_completion/hashcat.sh | 8 +- include/common.h | 8 +- include/types.h | 38 +++++--- src/autotune.c | 17 +++- src/combinator.c | 37 +++++--- src/dispatch.c | 38 +++----- src/interface.c | 159 ++++++++++++++++++++------------ src/opencl.c | 145 ++++++++++++++++++++--------- src/outfile.c | 25 +++-- src/selftest.c | 39 +++++--- src/straight.c | 4 + src/usage.c | 2 + src/user_options.c | 14 ++- src/wordlist.c | 111 ++++++++-------------- 17 files changed, 525 insertions(+), 263 deletions(-) create mode 100644 OpenCL/inc_scalar.cl diff --git a/OpenCL/inc_scalar.cl b/OpenCL/inc_scalar.cl new file mode 100644 index 000000000..bc29b2322 --- /dev/null +++ b/OpenCL/inc_scalar.cl @@ -0,0 +1,42 @@ +/** + * Author......: See docs/credits.txt + * License.....: MIT + */ + +#define COMPARE_S_SCALAR(h0,h1,h2,h3) \ +{ \ + if (((h0) == search[0]) && ((h1) == search[1]) && ((h2) == search[2]) && ((h3) == search[3])) \ + { \ + const u32 final_hash_pos = digests_offset + 0; \ + \ + if (atomic_inc (&hashes_shown[final_hash_pos]) == 0) \ + { \ + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, 0, final_hash_pos, gid, il_pos); \ + } \ + } \ +} + +#define COMPARE_M_SCALAR(h0,h1,h2,h3) \ +{ \ + const u32 digest_tp0[4] = { h0, h1, h2, h3 }; \ + \ + if (check (digest_tp0, \ + bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, \ + bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, \ + bitmap_mask, \ + bitmap_shift1, \ + bitmap_shift2)) \ + { \ + int digest_pos = find_hash (digest_tp0, digests_cnt, &digests_buf[digests_offset]); \ + \ + if (digest_pos != -1) \ + { \ + const u32 final_hash_pos = digests_offset + digest_pos; \ + \ + if (atomic_inc (&hashes_shown[final_hash_pos]) == 0) \ + { \ + mark_hash (plains_buf, d_return_buf, salt_pos, digests_cnt, digest_pos, final_hash_pos, gid, il_pos); \ + } \ + } \ + } \ +} diff --git a/OpenCL/m00000_a1.cl b/OpenCL/m00000_a1.cl index 4cfbccaf2..f65c9f20c 100644 --- a/OpenCL/m00000_a1.cl +++ b/OpenCL/m00000_a1.cl @@ -10,7 +10,9 @@ #include "inc_hash_functions.cl" #include "inc_types.cl" #include "inc_common.cl" +#include "inc_scalar.cl" #include "inc_simd.cl" +#include "inc_hash_md5.cl" __kernel void m00000_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { @@ -204,6 +206,48 @@ __kernel void m00000_m16 (__global pw_t *pws, __global const kernel_rule_t *rule { } +__kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +{ + /** + * modifier + */ + + const u32 lid = get_local_id (0); + const u32 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * base + */ + + md5_ctx_t ctx_outer; + + md5_init (&ctx_outer); + + md5_update_global (&ctx_outer, pws[gid].i, pws[gid].pw_len); + + /** + * loop + */ + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) + { + md5_ctx_t ctx_inner = ctx_outer; + + md5_update_global (&ctx_inner, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); + + md5_final (&ctx_inner); + + const u32 a = ctx_inner.h[0]; + const u32 b = ctx_inner.h[1]; + const u32 c = ctx_inner.h[2]; + const u32 d = ctx_inner.h[3]; + + COMPARE_M_SCALAR (a, d, c, b); + } +} + __kernel void m00000_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** @@ -410,3 +454,57 @@ __kernel void m00000_s08 (__global pw_t *pws, __global const kernel_rule_t *rule __kernel void m00000_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { } + +__kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +{ + /** + * modifier + */ + + const u32 lid = get_local_id (0); + const u32 gid = get_global_id (0); + + if (gid >= gid_max) return; + + /** + * digest + */ + + const u32 search[4] = + { + digests_buf[digests_offset].digest_buf[DGST_R0], + digests_buf[digests_offset].digest_buf[DGST_R1], + digests_buf[digests_offset].digest_buf[DGST_R2], + digests_buf[digests_offset].digest_buf[DGST_R3] + }; + + /** + * base + */ + + md5_ctx_t ctx_outer; + + md5_init (&ctx_outer); + + md5_update_global (&ctx_outer, pws[gid].i, pws[gid].pw_len); + + /** + * loop + */ + + for (u32 il_pos = 0; il_pos < il_cnt; il_pos++) + { + md5_ctx_t ctx_inner = ctx_outer; + + md5_update_global (&ctx_inner, combs_buf[il_pos].i, combs_buf[il_pos].pw_len); + + md5_final (&ctx_inner); + + const u32 a = ctx_inner.h[0]; + const u32 b = ctx_inner.h[1]; + const u32 c = ctx_inner.h[2]; + const u32 d = ctx_inner.h[3]; + + COMPARE_S_SCALAR (a, d, c, b); + } +} diff --git a/docs/changes.txt b/docs/changes.txt index 292be6af6..5795a73ea 100644 --- a/docs/changes.txt +++ b/docs/changes.txt @@ -4,7 +4,10 @@ ## Features ## +- Added more preparations to support to crack passwords and salts up to length 256 (unfinished yet) +- Added option --length-limit-disable to disable optimization based on password- and salt-length - Added self-test functionality for OpenCL kernels on startup +- Added option --self-test-disable to disable self-test functionality on startup ## ## Bugs diff --git a/extra/tab_completion/hashcat.sh b/extra/tab_completion/hashcat.sh index 25d11d6f9..f805fbcac 100644 --- a/extra/tab_completion/hashcat.sh +++ b/extra/tab_completion/hashcat.sh @@ -188,9 +188,9 @@ _hashcat () local HIDDEN_FILES_AGGRESIVE="exe|bin|pot|hcstat|dictstat|hcmask|hcchr|accepted|sh|cmd|restore" local BUILD_IN_CHARSETS='?l ?u ?d ?a ?b ?s ?h ?H' - local SHORT_OPTS="-m -a -V -v -h -b -t -o -p -c -d -w -n -u -j -k -r -g -1 -2 -3 -4 -i -I -s -l" - local LONG_OPTS="--hash-type --attack-mode --version --help --quiet --benchmark --hex-salt --hex-wordlist --hex-charset --force --status --status-timer --machine-readable --loopback --weak-hash-threshold --markov-hcstat --markov-disable --markov-classic --markov-threshold --runtime --session --speed-only --progress-only --restore --restore-file-path --restore-disable --outfile --outfile-format --outfile-autohex-disable --outfile-check-timer --outfile-check-dir --separator --show --left --username --remove --remove-timer --potfile-disable --potfile-path --debug-mode --debug-file --induction-dir --segment-size --bitmap-min --bitmap-max --cpu-affinity --opencl-info --opencl-devices --opencl-platforms --opencl-device-types --opencl-vector-width --workload-profile --kernel-accel --kernel-loops --nvidia-spin-damp --gpu-temp-disable --gpu-temp-abort --gpu-temp-retain --powertune-enable --skip --limit --keyspace --rule-left --rule-right --rules-file --generate-rules --generate-rules-func-min --generate-rules-func-max --generate-rules-seed --custom-charset1 --custom-charset2 --custom-charset3 --custom-charset4 --increment --increment-min --increment-max --logfile-disable --scrypt-tmto --truecrypt-keyfiles --stdout --keep-guessing --hccapx-message-pair --nonce-error-corrections --encoding-from --encoding-to" - local OPTIONS="-m -a -t -o -p -c -d -w -n -u -j -k -r -g -1 -2 -3 -4 -s -l --hash-type --attack-mode --status-timer --weak-hash-threshold --markov-hcstat --markov-threshold --runtime --session --timer --outfile --outfile-format --outfile-check-timer --outfile-check-dir --separator --remove-timer --potfile-path --restore-file-path --debug-mode --debug-file --induction-dir --segment-size --bitmap-min --bitmap-max --cpu-affinity --opencl-devices --opencl-platforms --opencl-device-types --opencl-vector-width --workload-profile --kernel-accel --kernel-loops --nvidia-spin-damp --gpu-temp-abort --gpu-temp-retain -disable --skip --limit --rule-left --rule-right --rules-file --generate-rules --generate-rules-func-min --generate-rules-func-max --generate-rules-seed --custom-charset1 --custom-charset2 --custom-charset3 --custom-charset4 --increment-min --increment-max --scrypt-tmto --truecrypt-keyfiles --hccapx-message-pair --nonce-error-corrections --encoding-from --encoding-to" + local SHORT_OPTS="-m -a -V -v -h -b -t -o -p -c -d -w -n -u -j -k -r -g -1 -2 -3 -4 -i -I -s -l -L" + local LONG_OPTS="--hash-type --attack-mode --version --help --quiet --benchmark --hex-salt --hex-wordlist --hex-charset --force --status --status-timer --machine-readable --loopback --weak-hash-threshold --markov-hcstat --markov-disable --markov-classic --markov-threshold --runtime --session --speed-only --progress-only --restore --restore-file-path --restore-disable --outfile --outfile-format --outfile-autohex-disable --outfile-check-timer --outfile-check-dir --separator --show --left --username --remove --remove-timer --potfile-disable --potfile-path --debug-mode --debug-file --induction-dir --segment-size --bitmap-min --bitmap-max --cpu-affinity --opencl-info --opencl-devices --opencl-platforms --opencl-device-types --opencl-vector-width --workload-profile --kernel-accel --kernel-loops --nvidia-spin-damp --gpu-temp-disable --gpu-temp-abort --gpu-temp-retain --powertune-enable --skip --limit --keyspace --rule-left --rule-right --rules-file --generate-rules --generate-rules-func-min --generate-rules-func-max --generate-rules-seed --custom-charset1 --custom-charset2 --custom-charset3 --custom-charset4 --increment --increment-min --increment-max --logfile-disable --scrypt-tmto --truecrypt-keyfiles --stdout --keep-guessing --hccapx-message-pair --nonce-error-corrections --encoding-from --encoding-to --length-limit-disable --self-test-disable" + local OPTIONS="-m -a -t -o -p -c -d -w -n -u -j -k -r -g -1 -2 -3 -4 -s -l --hash-type --attack-mode --status-timer --weak-hash-threshold --markov-hcstat --markov-threshold --runtime --session --timer --outfile --outfile-format --outfile-check-timer --outfile-check-dir --separator --remove-timer --potfile-path --restore-file-path --debug-mode --debug-file --induction-dir --segment-size --bitmap-min --bitmap-max --cpu-affinity --opencl-devices --opencl-platforms --opencl-device-types --opencl-vector-width --workload-profile --kernel-accel --kernel-loops --nvidia-spin-damp --gpu-temp-abort --gpu-temp-retain -disable --skip --limit --rule-left --rule-right --rules-file --generate-rules --generate-rules-func-min --generate-rules-func-max --generate-rules-seed --custom-charset1 --custom-charset2 --custom-charset3 --custom-charset4 --increment-min --increment-max --scrypt-tmto --truecrypt-keyfiles --hccapx-message-pair --nonce-error-corrections --encoding-from --encoding-to --length-limit-disable --self-test-disable" COMPREPLY=() local cur="${COMP_WORDS[COMP_CWORD]}" @@ -363,7 +363,7 @@ _hashcat () --status-timer|--markov-threshold|--runtime|--session|--separator|--segment-size|--rule-left|--rule-right| \ --nvidia-spin-damp|--gpu-temp-abort|--gpu-temp-retain|--generate-rules|--generate-rules-func-min|--generate-rules-func-max| \ --increment-min|--increment-max|--remove-timer|--bitmap-min|--bitmap-max|--skip|--limit|--generate-rules-seed| \ - --weak-hash-threshold|--outfile-check-timer|--outfile-check-dir|--induction-dir|--scrypt-tmto|--encoding-from|--encoding-to) + --weak-hash-threshold|--outfile-check-timer|--outfile-check-dir|--induction-dir|--scrypt-tmto|--encoding-from|--encoding-to|--length-limit-disable|--self-test-disable) return 0 ;; diff --git a/include/common.h b/include/common.h index 142e23dfa..6e5896f85 100644 --- a/include/common.h +++ b/include/common.h @@ -77,10 +77,8 @@ but this is nededed for VS compiler which doesn't have inline keyword but has __ // do not try to simply change this, it will not work #define PW_MIN 0 -#define PW_MAX 255 -#define PW_MAX1 (PW_MAX + 1) -#define PW_DICTMAX PW_MAX -#define PW_DICTMAX1 PW_MAX1 +#define PW_MAX 256 +#define PW_MAX_OLD 55 #define HCBUFSIZ_TINY 0x1000 #define HCBUFSIZ_LARGE 0x50000 @@ -91,7 +89,7 @@ but this is nededed for VS compiler which doesn't have inline keyword but has __ #define EXEC_CACHE 128 #define SPEED_CACHE 128 #define SPEED_MAXAGE 4096 -#define BLOCK_SIZE PW_MAX1 +#define BLOCK_SIZE 64 #define EXPECTED_ITERATIONS 10000 #if defined (_WIN) diff --git a/include/types.h b/include/types.h index 6da2f014c..f51c18830 100644 --- a/include/types.h +++ b/include/types.h @@ -236,8 +236,9 @@ typedef enum kern_run KERN_RUN_2 = 2000, KERN_RUN_23 = 2500, KERN_RUN_3 = 3000, - KERN_RUN_INIT2 = 4000, - KERN_RUN_LOOP2 = 5000 + KERN_RUN_4 = 4000, + KERN_RUN_INIT2 = 5000, + KERN_RUN_LOOP2 = 6000 } kern_run_t; @@ -531,6 +532,7 @@ typedef enum user_options_defaults KERNEL_LOOPS = 0, KEYSPACE = false, LEFT = false, + LENGTH_LIMIT_DISABLE = false, LIMIT = 0, LOGFILE_DISABLE = false, LOOPBACK = false, @@ -558,6 +560,7 @@ typedef enum user_options_defaults RP_GEN_SEED = 0, RUNTIME = 0, SCRYPT_TMTO = 0, + SELF_TEST_DISABLE = false, SEGMENT_SIZE = 33554432, SEPARATOR = ':', SHOW = false, @@ -610,6 +613,7 @@ typedef enum user_options_map IDX_KERNEL_LOOPS = 'u', IDX_KEYSPACE = 0xff14, IDX_LEFT = 0xff15, + IDX_LENGTH_LIMIT_DISABLE = 'L', IDX_LIMIT = 'l', IDX_LOGFILE_DISABLE = 0xff16, IDX_LOOPBACK = 0xff17, @@ -648,23 +652,24 @@ typedef enum user_options_map IDX_RULE_BUF_R = 'k', IDX_RUNTIME = 0xff30, IDX_SCRYPT_TMTO = 0xff31, + IDX_SELF_TEST_DISABLE = 0xff32, IDX_SEGMENT_SIZE = 'c', IDX_SEPARATOR = 'p', - IDX_SESSION = 0xff32, - IDX_SHOW = 0xff33, + IDX_SESSION = 0xff33, + IDX_SHOW = 0xff34, IDX_SKIP = 's', - IDX_STATUS = 0xff34, - IDX_STATUS_TIMER = 0xff35, - IDX_STDOUT_FLAG = 0xff36, - IDX_SPEED_ONLY = 0xff37, - IDX_PROGRESS_ONLY = 0xff38, - IDX_TRUECRYPT_KEYFILES = 0xff39, - IDX_USERNAME = 0xff3a, - IDX_VERACRYPT_KEYFILES = 0xff3b, - IDX_VERACRYPT_PIM = 0xff3c, + IDX_STATUS = 0xff35, + IDX_STATUS_TIMER = 0xff36, + IDX_STDOUT_FLAG = 0xff37, + IDX_SPEED_ONLY = 0xff38, + IDX_PROGRESS_ONLY = 0xff39, + IDX_TRUECRYPT_KEYFILES = 0xff3a, + IDX_USERNAME = 0xff3b, + IDX_VERACRYPT_KEYFILES = 0xff3c, + IDX_VERACRYPT_PIM = 0xff3d, IDX_VERSION_LOWER = 'v', IDX_VERSION = 'V', - IDX_WEAK_HASH_THRESHOLD = 0xff3d, + IDX_WEAK_HASH_THRESHOLD = 0xff3e, IDX_WORKLOAD_PROFILE = 'w' } user_options_map_t; @@ -915,6 +920,7 @@ typedef struct hc_device_param u32 kernel_threads_by_wgs_kernel2; u32 kernel_threads_by_wgs_kernel23; u32 kernel_threads_by_wgs_kernel3; + u32 kernel_threads_by_wgs_kernel4; u32 kernel_threads_by_wgs_kernel_init2; u32 kernel_threads_by_wgs_kernel_loop2; u32 kernel_threads_by_wgs_kernel_mp; @@ -979,6 +985,7 @@ typedef struct hc_device_param double exec_us_prev1[EXPECTED_ITERATIONS]; double exec_us_prev2[EXPECTED_ITERATIONS]; double exec_us_prev3[EXPECTED_ITERATIONS]; + double exec_us_prev4[EXPECTED_ITERATIONS]; double exec_us_prev_init2[EXPECTED_ITERATIONS]; double exec_us_prev_loop2[EXPECTED_ITERATIONS]; @@ -1011,6 +1018,7 @@ typedef struct hc_device_param cl_kernel kernel2; cl_kernel kernel23; cl_kernel kernel3; + cl_kernel kernel4; cl_kernel kernel_init2; cl_kernel kernel_loop2; cl_kernel kernel_mp; @@ -1446,6 +1454,7 @@ typedef struct user_options bool keep_guessing; bool keyspace; bool left; + bool length_limit_disable; bool logfile_disable; bool loopback; bool machine_readable; @@ -1459,6 +1468,7 @@ typedef struct user_options bool remove; bool restore; bool restore_disable; + bool self_test_disable; bool show; bool status; bool stdout_flag; diff --git a/src/autotune.c b/src/autotune.c index 0aa48998a..2d324fa10 100644 --- a/src/autotune.c +++ b/src/autotune.c @@ -13,7 +13,8 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const u32 kernel_accel, const u32 kernel_loops) { - hashconfig_t *hashconfig = hashcat_ctx->hashconfig; + hashconfig_t *hashconfig = hashcat_ctx->hashconfig; + user_options_t *user_options = hashcat_ctx->user_options; device_param->kernel_params_buf32[28] = 0; device_param->kernel_params_buf32[29] = kernel_loops; // not a bug, both need to be set @@ -21,9 +22,18 @@ static double try_run (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { - const u32 kernel_power_try = device_param->device_processors * device_param->kernel_threads_by_wgs_kernel1 * kernel_accel; + if (user_options->length_limit_disable == true) + { + const u32 kernel_power_try = device_param->device_processors * device_param->kernel_threads_by_wgs_kernel4 * kernel_accel; + + run_kernel (hashcat_ctx, device_param, KERN_RUN_4, kernel_power_try, true, 0); + } + else + { + const u32 kernel_power_try = device_param->device_processors * device_param->kernel_threads_by_wgs_kernel1 * kernel_accel; - run_kernel (hashcat_ctx, device_param, KERN_RUN_1, kernel_power_try, true, 0); + run_kernel (hashcat_ctx, device_param, KERN_RUN_1, kernel_power_try, true, 0); + } } else { @@ -291,6 +301,7 @@ static int autotune (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param memset (device_param->exec_us_prev1, 0, EXPECTED_ITERATIONS * sizeof (double)); memset (device_param->exec_us_prev2, 0, EXPECTED_ITERATIONS * sizeof (double)); memset (device_param->exec_us_prev3, 0, EXPECTED_ITERATIONS * sizeof (double)); + memset (device_param->exec_us_prev4, 0, EXPECTED_ITERATIONS * sizeof (double)); memset (device_param->exec_us_prev_init2, 0, EXPECTED_ITERATIONS * sizeof (double)); memset (device_param->exec_us_prev_loop2, 0, EXPECTED_ITERATIONS * sizeof (double)); diff --git a/src/combinator.c b/src/combinator.c index e3dea439b..d78527724 100644 --- a/src/combinator.c +++ b/src/combinator.c @@ -138,27 +138,35 @@ int combinator_ctx_init (hashcat_ctx_t *hashcat_ctx) combinator_ctx->dict1 = dictfile1; combinator_ctx->dict2 = dictfile2; - if (words1_cnt >= words2_cnt) + if (user_options->length_limit_disable == true) { combinator_ctx->combs_mode = COMBINATOR_MODE_BASE_LEFT; combinator_ctx->combs_cnt = words2_cnt; } else { - combinator_ctx->combs_mode = COMBINATOR_MODE_BASE_RIGHT; - combinator_ctx->combs_cnt = words1_cnt; + if (words1_cnt >= words2_cnt) + { + combinator_ctx->combs_mode = COMBINATOR_MODE_BASE_LEFT; + combinator_ctx->combs_cnt = words2_cnt; + } + else + { + combinator_ctx->combs_mode = COMBINATOR_MODE_BASE_RIGHT; + combinator_ctx->combs_cnt = words1_cnt; - // we also have to switch wordlist related rules! + // we also have to switch wordlist related rules! - char *tmpc = user_options->rule_buf_l; + char *tmpc = user_options->rule_buf_l; - user_options->rule_buf_l = user_options->rule_buf_r; - user_options->rule_buf_r = tmpc; + user_options->rule_buf_l = user_options->rule_buf_r; + user_options->rule_buf_r = tmpc; - u32 tmpi = user_options_extra->rule_len_l; + u32 tmpi = user_options_extra->rule_len_l; - user_options_extra->rule_len_l = user_options_extra->rule_len_r; - user_options_extra->rule_len_r = tmpi; + user_options_extra->rule_len_l = user_options_extra->rule_len_r; + user_options_extra->rule_len_r = tmpi; + } } } else if (user_options->attack_mode == ATTACK_MODE_BF) @@ -171,7 +179,14 @@ int combinator_ctx_init (hashcat_ctx_t *hashcat_ctx) } else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { - combinator_ctx->combs_mode = COMBINATOR_MODE_BASE_RIGHT; + if (user_options->length_limit_disable == true) + { + combinator_ctx->combs_mode = COMBINATOR_MODE_BASE_LEFT; + } + else + { + combinator_ctx->combs_mode = COMBINATOR_MODE_BASE_RIGHT; + } } return 0; diff --git a/src/dispatch.c b/src/dispatch.c index 71bd1f871..73dc57867 100644 --- a/src/dispatch.c +++ b/src/dispatch.c @@ -172,15 +172,10 @@ static int calc_stdin (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par const size_t iconv_rc = iconv (iconv_ctx, &line_buf, &line_len, &iconv_ptr, &iconv_sz); - if (iconv_rc == (size_t) -1) - { - line_len = PW_MAX1; - } - else - { - line_buf = iconv_tmp; - line_len = HCBUFSIZ_TINY - iconv_sz; - } + if (iconv_rc == (size_t) -1) continue; + + line_buf = iconv_tmp; + line_len = HCBUFSIZ_TINY - iconv_sz; } // post-process rule engine @@ -189,25 +184,19 @@ static int calc_stdin (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_par if (run_rule_engine ((int) user_options_extra->rule_len_l, user_options->rule_buf_l)) { + if (line_len >= BLOCK_SIZE) continue; + memset (rule_buf_out, 0, sizeof (rule_buf_out)); - int rule_len_out = -1; - - if (line_len < BLOCK_SIZE) - { - rule_len_out = _old_apply_rule (user_options->rule_buf_l, (int) user_options_extra->rule_len_l, line_buf, (int) line_len, rule_buf_out); - } + const int rule_len_out = _old_apply_rule (user_options->rule_buf_l, (int) user_options_extra->rule_len_l, line_buf, (int) line_len, rule_buf_out); if (rule_len_out < 0) continue; line_buf = rule_buf_out; - line_len = (u32) rule_len_out; + line_len = (size_t) rule_len_out; } - if (line_len > PW_MAX) - { - continue; - } + if (line_len >= PW_MAX) continue; // hmm that's always the case, or? @@ -505,14 +494,11 @@ static int calc (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param) if (run_rule_engine ((int) user_options_extra->rule_len_l, user_options->rule_buf_l)) { - memset (rule_buf_out, 0, sizeof (rule_buf_out)); + if (line_len >= BLOCK_SIZE) continue; - int rule_len_out = -1; + memset (rule_buf_out, 0, sizeof (rule_buf_out)); - if (line_len < BLOCK_SIZE) - { - rule_len_out = _old_apply_rule (user_options->rule_buf_l, (int) user_options_extra->rule_len_l, line_buf, (int) line_len, rule_buf_out); - } + const int rule_len_out = _old_apply_rule (user_options->rule_buf_l, (int) user_options_extra->rule_len_l, line_buf, (int) line_len, rule_buf_out); if (rule_len_out < 0) continue; diff --git a/src/interface.c b/src/interface.c index 1c777425b..68990152c 100644 --- a/src/interface.c +++ b/src/interface.c @@ -24276,6 +24276,30 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) hashconfig->opts_type |= OPTS_TYPE_PT_NEVERCRACK; } + if (user_options->length_limit_disable == true) + { + hashconfig->opts_type &= ~OPTS_TYPE_PT_UTF16LE; + hashconfig->opts_type &= ~OPTS_TYPE_PT_UTF16BE; + hashconfig->opts_type &= ~OPTS_TYPE_PT_ADD01; + hashconfig->opts_type &= ~OPTS_TYPE_PT_ADD02; + hashconfig->opts_type &= ~OPTS_TYPE_PT_ADD80; + hashconfig->opts_type &= ~OPTS_TYPE_PT_ADDBITS14; + hashconfig->opts_type &= ~OPTS_TYPE_PT_ADDBITS15; + hashconfig->opts_type &= ~OPTS_TYPE_ST_UTF16LE; + hashconfig->opts_type &= ~OPTS_TYPE_ST_UTF16BE; + hashconfig->opts_type &= ~OPTS_TYPE_ST_ADD01; + hashconfig->opts_type &= ~OPTS_TYPE_ST_ADD02; + hashconfig->opts_type &= ~OPTS_TYPE_ST_ADD80; + hashconfig->opts_type &= ~OPTS_TYPE_ST_ADDBITS14; + hashconfig->opts_type &= ~OPTS_TYPE_ST_ADDBITS15; + + hashconfig->opti_type &= ~OPTI_TYPE_PRECOMPUTE_INIT; + hashconfig->opti_type &= ~OPTI_TYPE_PRECOMPUTE_MERKLE; + hashconfig->opti_type &= ~OPTI_TYPE_MEET_IN_MIDDLE; + hashconfig->opti_type &= ~OPTI_TYPE_PREPENDED_SALT; + hashconfig->opti_type &= ~OPTI_TYPE_APPENDED_SALT; + } + const u32 is_salted = ((hashconfig->salt_type == SALT_TYPE_INTERN) | (hashconfig->salt_type == SALT_TYPE_EXTERN) | (hashconfig->salt_type == SALT_TYPE_EMBEDDED) @@ -24511,91 +24535,104 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx) // pw_max - hashconfig->pw_max = PW_MAX; - - if ((hashconfig->opts_type & OPTS_TYPE_PT_UTF16LE) || (hashconfig->opts_type & OPTS_TYPE_PT_UTF16BE)) + if (user_options->length_limit_disable == false) + { + hashconfig->pw_max = PW_MAX; + } + else { - hashconfig->pw_max = PW_MAX / 2; + hashconfig->pw_max = PW_MAX_OLD; + + if ((hashconfig->opts_type & OPTS_TYPE_PT_UTF16LE) || (hashconfig->opts_type & OPTS_TYPE_PT_UTF16BE)) + { + hashconfig->pw_max /= 2; + } + + switch (hashconfig->hash_mode) + { + case 125: hashconfig->pw_max = 32; + break; + case 500: hashconfig->pw_max = 16; + break; + case 1600: hashconfig->pw_max = 16; + break; + case 1800: hashconfig->pw_max = 16; + break; + case 2100: hashconfig->pw_max = 27; + break; + case 5200: hashconfig->pw_max = 24; + break; + case 5800: hashconfig->pw_max = 16; + break; + case 6300: hashconfig->pw_max = 16; + break; + case 7000: hashconfig->pw_max = 19; + break; + case 7400: hashconfig->pw_max = 16; + break; + case 7700: hashconfig->pw_max = 8; + break; + case 7900: hashconfig->pw_max = 48; + break; + case 8600: hashconfig->pw_max = 16; + break; + case 10300: hashconfig->pw_max = 40; + break; + case 10500: hashconfig->pw_max = 40; + break; + case 10700: hashconfig->pw_max = 16; + break; + case 11300: hashconfig->pw_max = 40; + break; + case 11600: hashconfig->pw_max = 32; + break; + case 12500: hashconfig->pw_max = 20; + break; + case 12800: hashconfig->pw_max = 24; + break; + case 14400: hashconfig->pw_max = 24; + break; + case 15400: hashconfig->pw_max = 32; + break; + case 15500: hashconfig->pw_max = 16; + break; + } + + // fully converted to length 256 + + switch (hashconfig->hash_mode) + { + case 400: hashconfig->pw_max = 256; + break; + case 2100: hashconfig->pw_max = 256; + break; + } } + // pw_max : algo specific hard limits + switch (hashconfig->hash_mode) { - case 125: hashconfig->pw_max = 32; - break; - case 400: hashconfig->pw_max = 40; - break; - case 500: hashconfig->pw_max = 16; - break; case 1500: hashconfig->pw_max = 8; break; - case 1600: hashconfig->pw_max = 16; - break; - case 1800: hashconfig->pw_max = 16; - break; - case 2100: hashconfig->pw_max = 27; + case 2500: hashconfig->pw_max = 64; break; case 3000: hashconfig->pw_max = 7; break; - case 5200: hashconfig->pw_max = 24; - break; - case 5800: hashconfig->pw_max = 16; - break; - case 6300: hashconfig->pw_max = 16; - break; - case 7000: hashconfig->pw_max = 19; - break; - case 7400: hashconfig->pw_max = 16; - break; - case 7700: hashconfig->pw_max = 8; - break; - case 7900: hashconfig->pw_max = 48; - break; case 8500: hashconfig->pw_max = 8; break; - case 8600: hashconfig->pw_max = 16; - break; case 9710: hashconfig->pw_max = 5; break; case 9810: hashconfig->pw_max = 5; break; case 10410: hashconfig->pw_max = 5; break; - case 10300: hashconfig->pw_max = 40; - break; - case 10500: hashconfig->pw_max = 40; - break; - case 10700: hashconfig->pw_max = 16; - break; - case 11300: hashconfig->pw_max = 40; - break; - case 11600: hashconfig->pw_max = 32; - break; - case 12500: hashconfig->pw_max = 20; - break; - case 12800: hashconfig->pw_max = 24; - break; case 14000: hashconfig->pw_max = 8; break; case 14100: hashconfig->pw_max = 24; break; - case 14400: hashconfig->pw_max = 24; - break; case 14900: hashconfig->pw_max = 10; break; - case 15400: hashconfig->pw_max = 32; - break; - case 15500: hashconfig->pw_max = 16; - break; - } - - // converted to length 256 - - switch (hashconfig->hash_mode) - { - case 400: hashconfig->pw_max = 256; - break; - case 2100: hashconfig->pw_max = 256; - break; } return 0; diff --git a/src/opencl.c b/src/opencl.c index d11a17caa..b0f6a55ac 100644 --- a/src/opencl.c +++ b/src/opencl.c @@ -1125,23 +1125,32 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, } } - if (highest_pw_len < 16) + if (user_options->length_limit_disable == true) { - CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_1, pws_cnt, true, fast_iteration); - - if (CL_rc == -1) return -1; - } - else if (highest_pw_len < 32) - { - CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_2, pws_cnt, true, fast_iteration); + CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_4, pws_cnt, true, fast_iteration); if (CL_rc == -1) return -1; } else { - CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_cnt, true, fast_iteration); + if (highest_pw_len < 16) + { + CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_1, pws_cnt, true, fast_iteration); - if (CL_rc == -1) return -1; + if (CL_rc == -1) return -1; + } + else if (highest_pw_len < 32) + { + CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_2, pws_cnt, true, fast_iteration); + + if (CL_rc == -1) return -1; + } + else + { + CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_cnt, true, fast_iteration); + + if (CL_rc == -1) return -1; + } } } else @@ -1358,6 +1367,10 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con kernel = device_param->kernel3; kernel_threads = device_param->kernel_threads_by_wgs_kernel3; break; + case KERN_RUN_4: + kernel = device_param->kernel4; + kernel_threads = device_param->kernel_threads_by_wgs_kernel4; + break; case KERN_RUN_INIT2: kernel = device_param->kernel_init2; kernel_threads = device_param->kernel_threads_by_wgs_kernel_init2; @@ -1449,6 +1462,7 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con case KERN_RUN_1: if (device_param->exec_us_prev1[iteration] > 0) usleep ((useconds_t)(device_param->exec_us_prev1[iteration] * device_param->nvidia_spin_damp)); break; case KERN_RUN_2: if (device_param->exec_us_prev2[iteration] > 0) usleep ((useconds_t)(device_param->exec_us_prev2[iteration] * device_param->nvidia_spin_damp)); break; case KERN_RUN_3: if (device_param->exec_us_prev3[iteration] > 0) usleep ((useconds_t)(device_param->exec_us_prev3[iteration] * device_param->nvidia_spin_damp)); break; + case KERN_RUN_4: if (device_param->exec_us_prev4[iteration] > 0) usleep ((useconds_t)(device_param->exec_us_prev4[iteration] * device_param->nvidia_spin_damp)); break; case KERN_RUN_INIT2: if (device_param->exec_us_prev_init2[iteration] > 0) usleep ((useconds_t)(device_param->exec_us_prev_init2[iteration] * device_param->nvidia_spin_damp)); break; case KERN_RUN_LOOP2: if (device_param->exec_us_prev_loop2[iteration] > 0) usleep ((useconds_t)(device_param->exec_us_prev_loop2[iteration] * device_param->nvidia_spin_damp)); break; } @@ -1477,6 +1491,7 @@ int run_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, con case KERN_RUN_1: device_param->exec_us_prev1[iteration] = exec_us; break; case KERN_RUN_2: device_param->exec_us_prev2[iteration] = exec_us; break; case KERN_RUN_3: device_param->exec_us_prev3[iteration] = exec_us; break; + case KERN_RUN_4: device_param->exec_us_prev4[iteration] = exec_us; break; case KERN_RUN_INIT2: device_param->exec_us_prev_init2[iteration] = exec_us; break; case KERN_RUN_LOOP2: device_param->exec_us_prev_loop2[iteration] = exec_us; break; } @@ -1739,9 +1754,42 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) { - if (user_options->attack_mode == ATTACK_MODE_COMBI) + if (user_options->length_limit_disable == true) + { + + + } + else { - if (combinator_ctx->combs_mode == COMBINATOR_MODE_BASE_RIGHT) + if (user_options->attack_mode == ATTACK_MODE_COMBI) + { + if (combinator_ctx->combs_mode == COMBINATOR_MODE_BASE_RIGHT) + { + if (hashconfig->opts_type & OPTS_TYPE_PT_ADD01) + { + for (u32 i = 0; i < pws_cnt; i++) + { + const u32 pw_len = device_param->pws_buf[i].pw_len; + + u8 *ptr = (u8 *) device_param->pws_buf[i].i; + + ptr[pw_len] = 0x01; + } + } + else if (hashconfig->opts_type & OPTS_TYPE_PT_ADD80) + { + for (u32 i = 0; i < pws_cnt; i++) + { + const u32 pw_len = device_param->pws_buf[i].pw_len; + + u8 *ptr = (u8 *) device_param->pws_buf[i].i; + + ptr[pw_len] = 0x80; + } + } + } + } + else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) { if (hashconfig->opts_type & OPTS_TYPE_PT_ADD01) { @@ -1766,38 +1814,13 @@ int run_copy (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, const } } } - } - else if (user_options->attack_mode == ATTACK_MODE_HYBRID2) - { - if (hashconfig->opts_type & OPTS_TYPE_PT_ADD01) - { - for (u32 i = 0; i < pws_cnt; i++) - { - const u32 pw_len = device_param->pws_buf[i].pw_len; - - u8 *ptr = (u8 *) device_param->pws_buf[i].i; - ptr[pw_len] = 0x01; - } - } - else if (hashconfig->opts_type & OPTS_TYPE_PT_ADD80) - { - for (u32 i = 0; i < pws_cnt; i++) - { - const u32 pw_len = device_param->pws_buf[i].pw_len; + int CL_rc; - u8 *ptr = (u8 *) device_param->pws_buf[i].i; + CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL); - ptr[pw_len] = 0x80; - } - } + if (CL_rc == -1) return -1; } - - int CL_rc; - - CL_rc = hc_clEnqueueWriteBuffer (hashcat_ctx, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, pws_cnt * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL); - - if (CL_rc == -1) return -1; } else if (user_options_extra->attack_kern == ATTACK_KERN_BF) { @@ -1945,19 +1968,21 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co int line_len = fgetl (combs_fp, line_buf); - if (line_len >= PW_MAX1) continue; - line_len = convert_from_hex (hashcat_ctx, line_buf, line_len); + if (line_len >= PW_MAX) continue; + char *line_buf_new = line_buf; char rule_buf_out[BLOCK_SIZE]; if (run_rule_engine (user_options_extra->rule_len_r, user_options->rule_buf_r)) { + if (line_len >= BLOCK_SIZE) continue; + memset (rule_buf_out, 0, sizeof (rule_buf_out)); - int rule_len_out = _old_apply_rule (user_options->rule_buf_r, user_options_extra->rule_len_r, line_buf, line_len, rule_buf_out); + const int rule_len_out = _old_apply_rule (user_options->rule_buf_r, user_options_extra->rule_len_r, line_buf, line_len, rule_buf_out); if (rule_len_out < 0) { @@ -1971,13 +1996,13 @@ int run_cracker (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, co line_buf_new = rule_buf_out; } - line_len = MIN (line_len, PW_DICTMAX); + line_len = MIN (line_len, PW_MAX - 1); u8 *ptr = (u8 *) device_param->combs_buf[i].i; memcpy (ptr, line_buf_new, line_len); - memset (ptr + line_len, 0, PW_DICTMAX1 - line_len); + memset (ptr + line_len, 0, PW_MAX - line_len); if (hashconfig->opts_type & OPTS_TYPE_PT_UPPER) { @@ -4800,6 +4825,15 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) CL_rc = hc_clCreateKernel (hashcat_ctx, device_param->program, kernel_name, &device_param->kernel3); if (CL_rc == -1) return -1; + + if (user_options->length_limit_disable == true) + { + snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05u_sxx", hashconfig->kern_type); + + CL_rc = hc_clCreateKernel (hashcat_ctx, device_param->program, kernel_name, &device_param->kernel4); + + if (CL_rc == -1) return -1; + } } else { @@ -4820,6 +4854,15 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) CL_rc = hc_clCreateKernel (hashcat_ctx, device_param->program, kernel_name, &device_param->kernel3); if (CL_rc == -1) return -1; + + if (user_options->length_limit_disable == true) + { + snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05u_mxx", hashconfig->kern_type); + + CL_rc = hc_clCreateKernel (hashcat_ctx, device_param->program, kernel_name, &device_param->kernel4); + + if (CL_rc == -1) return -1; + } } if (user_options->attack_mode == ATTACK_MODE_BF) @@ -4943,6 +4986,15 @@ int opencl_session_begin (hashcat_ctx_t *hashcat_ctx) if (CL_rc == -1) return -1; + // kernel4 + + if (user_options->length_limit_disable == true) + { + CL_rc = get_kernel_threads (hashcat_ctx, device_param, device_param->kernel4, &device_param->kernel_threads_by_wgs_kernel4); + + if (CL_rc == -1) return -1; + } + // GPU memset CL_rc = hc_clCreateKernel (hashcat_ctx, device_param->program, "gpu_memset", &device_param->kernel_memset); @@ -5209,6 +5261,7 @@ void opencl_session_destroy (hashcat_ctx_t *hashcat_ctx) if (device_param->kernel2) hc_clReleaseKernel (hashcat_ctx, device_param->kernel2); if (device_param->kernel23) hc_clReleaseKernel (hashcat_ctx, device_param->kernel23); if (device_param->kernel3) hc_clReleaseKernel (hashcat_ctx, device_param->kernel3); + if (device_param->kernel4) hc_clReleaseKernel (hashcat_ctx, device_param->kernel4); if (device_param->kernel_init2) hc_clReleaseKernel (hashcat_ctx, device_param->kernel_init2); if (device_param->kernel_loop2) hc_clReleaseKernel (hashcat_ctx, device_param->kernel_loop2); if (device_param->kernel_mp) hc_clReleaseKernel (hashcat_ctx, device_param->kernel_mp); @@ -5269,6 +5322,7 @@ void opencl_session_destroy (hashcat_ctx_t *hashcat_ctx) device_param->kernel2 = NULL; device_param->kernel23 = NULL; device_param->kernel3 = NULL; + device_param->kernel4 = NULL; device_param->kernel_init2 = NULL; device_param->kernel_loop2 = NULL; device_param->kernel_mp = NULL; @@ -5349,6 +5403,7 @@ int opencl_session_update_combinator (hashcat_ctx_t *hashcat_ctx) CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->kernel1, 33, sizeof (cl_uint), device_param->kernel_params[33]); if (CL_rc == -1) return -1; CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->kernel2, 33, sizeof (cl_uint), device_param->kernel_params[33]); if (CL_rc == -1) return -1; CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->kernel3, 33, sizeof (cl_uint), device_param->kernel_params[33]); if (CL_rc == -1) return -1; + CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->kernel4, 33, sizeof (cl_uint), device_param->kernel_params[33]); if (CL_rc == -1) return -1; if (hashconfig->opts_type & OPTS_TYPE_HOOK12) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->kernel12, 33, sizeof (cl_uint), device_param->kernel_params[33]); if (CL_rc == -1) return -1; } if (hashconfig->opts_type & OPTS_TYPE_HOOK23) { CL_rc = hc_clSetKernelArg (hashcat_ctx, device_param->kernel23, 33, sizeof (cl_uint), device_param->kernel_params[33]); if (CL_rc == -1) return -1; } diff --git a/src/outfile.c b/src/outfile.c index 94843d676..aefea3bdc 100644 --- a/src/outfile.c +++ b/src/outfile.c @@ -84,20 +84,29 @@ int build_plain (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param, pl memcpy (plain_ptr, comb_buf, comb_len); } - int pw_max_combi; + plain_len += comb_len; - if (hashconfig->pw_max < PW_DICTMAX) + if (user_options->length_limit_disable == true) { - pw_max_combi = hashconfig->pw_max; + int pw_max_combi; + + #define PW_DICTMAX 32 + + if (hashconfig->pw_max < PW_DICTMAX) + { + pw_max_combi = hashconfig->pw_max; + } + else + { + pw_max_combi = PW_MAX_OLD; + } + + plain_len = MIN ((int) plain_len, (int) pw_max_combi); } else { - pw_max_combi = PW_MAX; + plain_len = MIN ((int) plain_len, (int) hashconfig->pw_max); } - - plain_len += comb_len; - - if (plain_len > pw_max_combi) plain_len = pw_max_combi; } else if (user_options->attack_mode == ATTACK_MODE_BF) { diff --git a/src/selftest.c b/src/selftest.c index c6da113d9..c6f9e9896 100644 --- a/src/selftest.c +++ b/src/selftest.c @@ -18,6 +18,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param hashconfig_t *hashconfig = hashcat_ctx->hashconfig; hashes_t *hashes = hashcat_ctx->hashes; status_ctx_t *status_ctx = hashcat_ctx->status_ctx; + user_options_t *user_options = hashcat_ctx->user_options; user_options_extra_t *user_options_extra = hashcat_ctx->user_options_extra; cl_int CL_err; @@ -40,6 +41,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param const u32 kernel_threads_by_wgs_kernel2_sav = device_param->kernel_threads_by_wgs_kernel2; const u32 kernel_threads_by_wgs_kernel23_sav = device_param->kernel_threads_by_wgs_kernel23; const u32 kernel_threads_by_wgs_kernel3_sav = device_param->kernel_threads_by_wgs_kernel3; + const u32 kernel_threads_by_wgs_kernel4_sav = device_param->kernel_threads_by_wgs_kernel4; const u32 kernel_threads_by_wgs_kernel_init2_sav = device_param->kernel_threads_by_wgs_kernel_init2; const u32 kernel_threads_by_wgs_kernel_loop2_sav = device_param->kernel_threads_by_wgs_kernel_loop2; @@ -58,6 +60,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (device_param->kernel_threads_by_user == device_param->kernel_threads_by_wgs_kernel2) device_param->kernel_threads_by_wgs_kernel2 = 1; if (device_param->kernel_threads_by_user == device_param->kernel_threads_by_wgs_kernel23) device_param->kernel_threads_by_wgs_kernel23 = 1; if (device_param->kernel_threads_by_user == device_param->kernel_threads_by_wgs_kernel3) device_param->kernel_threads_by_wgs_kernel3 = 1; + if (device_param->kernel_threads_by_user == device_param->kernel_threads_by_wgs_kernel4) device_param->kernel_threads_by_wgs_kernel4 = 1; if (device_param->kernel_threads_by_user == device_param->kernel_threads_by_wgs_kernel_init2) device_param->kernel_threads_by_wgs_kernel_init2 = 1; if (device_param->kernel_threads_by_user == device_param->kernel_threads_by_wgs_kernel_loop2) device_param->kernel_threads_by_wgs_kernel_loop2 = 1; } @@ -303,23 +306,32 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param if (hashconfig->attack_exec == ATTACK_EXEC_INSIDE_KERNEL) { - if (highest_pw_len < 16) + if (user_options->length_limit_disable == true) { - CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 1, false, 0); - - if (CL_rc == -1) return -1; - } - else if (highest_pw_len < 32) - { - CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 1, false, 0); + CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_4, 1, false, 0); if (CL_rc == -1) return -1; } else { - CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_3, 1, false, 0); + if (highest_pw_len < 16) + { + CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_1, 1, false, 0); - if (CL_rc == -1) return -1; + if (CL_rc == -1) return -1; + } + else if (highest_pw_len < 32) + { + CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_2, 1, false, 0); + + if (CL_rc == -1) return -1; + } + else + { + CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_3, 1, false, 0); + + if (CL_rc == -1) return -1; + } } } else @@ -426,6 +438,7 @@ static int selftest (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param device_param->kernel_threads_by_wgs_kernel2 = kernel_threads_by_wgs_kernel2_sav; device_param->kernel_threads_by_wgs_kernel23 = kernel_threads_by_wgs_kernel23_sav; device_param->kernel_threads_by_wgs_kernel3 = kernel_threads_by_wgs_kernel3_sav; + device_param->kernel_threads_by_wgs_kernel4 = kernel_threads_by_wgs_kernel4_sav; device_param->kernel_threads_by_wgs_kernel_init2 = kernel_threads_by_wgs_kernel_init2_sav; device_param->kernel_threads_by_wgs_kernel_loop2 = kernel_threads_by_wgs_kernel_loop2_sav; @@ -499,6 +512,10 @@ void *thread_selftest (void *p) if (opencl_ctx->enabled == false) return NULL; + user_options_t *user_options = hashcat_ctx->user_options; + + if (user_options->self_test_disable == true) return NULL; + hc_device_param_t *device_param = opencl_ctx->devices_param + thread_param->tid; if (device_param->skipped == true) return NULL; @@ -507,7 +524,7 @@ void *thread_selftest (void *p) if (rc_selftest == -1) { - // we should do something here, tell hashcat main that autotune failed to abort + // we should do something here, tell hashcat main that selftest failed to abort } return NULL; diff --git a/src/straight.c b/src/straight.c index 7b7d43001..fe805636f 100644 --- a/src/straight.c +++ b/src/straight.c @@ -253,6 +253,10 @@ int straight_ctx_init (hashcat_ctx_t *hashcat_ctx) const bool has_noop = kernel_rules_has_noop (straight_ctx->kernel_rules_buf, straight_ctx->kernel_rules_cnt); + // this entire section should go away as soon as we have a rule engine for 256 byte + + #define PW_DICTMAX 32 + if (has_noop == false) { switch (user_options_extra->attack_kern) diff --git a/src/usage.c b/src/usage.c index 4b2664a60..c71ce51ba 100644 --- a/src/usage.c +++ b/src/usage.c @@ -37,6 +37,7 @@ static const char *USAGE_BIG[] = " --status-timer | Num | Sets seconds between status screen updates to X | --status-timer=1", " --machine-readable | | Display the status view in a machine-readable format |", " --keep-guessing | | Keep guessing the hash after it has been cracked |", + " --self-test-disable | | Disable self-test functionality on startup |", " --loopback | | Add new plains to induct directory |", " --weak-hash-threshold | Num | Threshold X when to stop checking for weak hashes | --weak=0", " --markov-hcstat | File | Specify hcstat file to use | --markov-hc=my.hcstat", @@ -76,6 +77,7 @@ static const char *USAGE_BIG[] = " -b, --benchmark | | Run benchmark |", " --speed-only | | Return expected speed of the attack, then quit |", " --progress-only | | Return ideal progress step size and time to process |", + " -L, --length-limit-disable | | Disable optimizations for password- and salt-length |", " -c, --segment-size | Num | Sets size in MB to cache from the wordfile to X | -c 32", " --bitmap-min | Num | Sets minimum bits allowed for bitmaps to X | --bitmap-min=24", " --bitmap-max | Num | Sets maximum bits allowed for bitmaps to X | --bitmap-max=24", diff --git a/src/user_options.c b/src/user_options.c index 17c5ddc6e..5dc5a20bd 100644 --- a/src/user_options.c +++ b/src/user_options.c @@ -14,7 +14,7 @@ #include "outfile.h" #include "user_options.h" -static const char short_options[] = "hVvm:a:r:j:k:g:o:t:d:D:n:u:c:p:s:l:1:2:3:4:iIbw:"; +static const char short_options[] = "hVvm:a:r:j:k:g:o:t:d:D:n:u:c:p:s:l:1:2:3:4:iIbw:L"; static const struct option long_options[] = { @@ -55,6 +55,7 @@ static const struct option long_options[] = {"kernel-loops", required_argument, 0, IDX_KERNEL_LOOPS}, {"keyspace", no_argument, 0, IDX_KEYSPACE}, {"left", no_argument, 0, IDX_LEFT}, + {"length-limit-disable", no_argument, 0, IDX_LENGTH_LIMIT_DISABLE}, {"limit", required_argument, 0, IDX_LIMIT}, {"logfile-disable", no_argument, 0, IDX_LOGFILE_DISABLE}, {"loopback", no_argument, 0, IDX_LOOPBACK}, @@ -89,6 +90,7 @@ static const struct option long_options[] = {"rules-file", required_argument, 0, IDX_RP_FILE}, {"runtime", required_argument, 0, IDX_RUNTIME}, {"scrypt-tmto", required_argument, 0, IDX_SCRYPT_TMTO}, + {"self-test-disable", no_argument, 0, IDX_SELF_TEST_DISABLE}, {"segment-size", required_argument, 0, IDX_SEGMENT_SIZE}, {"separator", required_argument, 0, IDX_SEPARATOR}, {"seperator", required_argument, 0, IDX_SEPARATOR}, @@ -156,6 +158,7 @@ int user_options_init (hashcat_ctx_t *hashcat_ctx) user_options->keep_guessing = KEEP_GUESSING; user_options->keyspace = KEYSPACE; user_options->left = LEFT; + user_options->length_limit_disable = LENGTH_LIMIT_DISABLE; user_options->limit = LIMIT; user_options->logfile_disable = LOGFILE_DISABLE; user_options->loopback = LOOPBACK; @@ -194,6 +197,7 @@ int user_options_init (hashcat_ctx_t *hashcat_ctx) user_options->rule_buf_r = RULE_BUF_R; user_options->runtime = RUNTIME; user_options->scrypt_tmto = SCRYPT_TMTO; + user_options->self_test_disable = SELF_TEST_DISABLE; user_options->segment_size = SEGMENT_SIZE; user_options->separator = SEPARATOR; user_options->session = PROGNAME; @@ -312,6 +316,7 @@ int user_options_getopt (hashcat_ctx_t *hashcat_ctx, int argc, char **argv) case IDX_QUIET: user_options->quiet = true; break; case IDX_SHOW: user_options->show = true; break; case IDX_LEFT: user_options->left = true; break; + case IDX_LENGTH_LIMIT_DISABLE: user_options->length_limit_disable = true; break; case IDX_ADVICE_DISABLE: user_options->advice_disable = true; break; case IDX_USERNAME: user_options->username = true; break; case IDX_REMOVE: user_options->remove = true; break; @@ -326,6 +331,7 @@ int user_options_getopt (hashcat_ctx_t *hashcat_ctx, int argc, char **argv) case IDX_INDUCTION_DIR: user_options->induction_dir = optarg; break; case IDX_OUTFILE_CHECK_DIR: user_options->outfile_check_dir = optarg; break; case IDX_FORCE: user_options->force = true; break; + case IDX_SELF_TEST_DISABLE: user_options->self_test_disable = true; break; case IDX_SKIP: user_options->skip = atoll (optarg); break; case IDX_LIMIT: user_options->limit = atoll (optarg); break; case IDX_KEEP_GUESSING: user_options->keep_guessing = true; break; @@ -2121,6 +2127,7 @@ void user_options_logger (hashcat_ctx_t *hashcat_ctx) logfile_top_uint (user_options->kernel_loops); logfile_top_uint (user_options->keyspace); logfile_top_uint (user_options->left); + logfile_top_uint (user_options->length_limit_disable); logfile_top_uint (user_options->logfile_disable); logfile_top_uint (user_options->loopback); logfile_top_uint (user_options->machine_readable); @@ -2135,6 +2142,7 @@ void user_options_logger (hashcat_ctx_t *hashcat_ctx) logfile_top_uint (user_options->outfile_format); logfile_top_uint (user_options->potfile_disable); logfile_top_uint (user_options->powertune_enable); + logfile_top_uint (user_options->progress_only); logfile_top_uint (user_options->quiet); logfile_top_uint (user_options->remove); logfile_top_uint (user_options->remove_timer); @@ -2149,12 +2157,12 @@ void user_options_logger (hashcat_ctx_t *hashcat_ctx) logfile_top_uint (user_options->runtime); logfile_top_uint (user_options->scrypt_tmto); logfile_top_uint (user_options->segment_size); + logfile_top_uint (user_options->self_test_disable); logfile_top_uint (user_options->show); + logfile_top_uint (user_options->speed_only); logfile_top_uint (user_options->status); logfile_top_uint (user_options->status_timer); logfile_top_uint (user_options->stdout_flag); - logfile_top_uint (user_options->speed_only); - logfile_top_uint (user_options->progress_only); logfile_top_uint (user_options->usage); logfile_top_uint (user_options->username); logfile_top_uint (user_options->veracrypt_pim); diff --git a/src/wordlist.c b/src/wordlist.c index f814bad23..2c02ec2f4 100644 --- a/src/wordlist.c +++ b/src/wordlist.c @@ -195,46 +195,27 @@ void get_next_word (hashcat_ctx_t *hashcat_ctx, FILE *fd, char **out_buf, u32 *o const size_t iconv_rc = iconv (wl_data->iconv_ctx, &ptr, &ptr_len, &iconv_ptr, &iconv_sz); - if (iconv_rc == (size_t) -1) - { - len = PW_MAX1; - } - else - { - ptr = wl_data->iconv_tmp; - len = HCBUFSIZ_TINY - iconv_sz; - } + if (iconv_rc == (size_t) -1) continue; + + ptr = wl_data->iconv_tmp; + len = HCBUFSIZ_TINY - iconv_sz; } if (run_rule_engine (user_options_extra->rule_len_l, user_options->rule_buf_l)) { - int rule_len_out = -1; + if (len >= BLOCK_SIZE) continue; - if (len < BLOCK_SIZE) - { - char unused[BLOCK_SIZE] = { 0 }; + char rule_buf_out[BLOCK_SIZE]; - rule_len_out = _old_apply_rule (user_options->rule_buf_l, user_options_extra->rule_len_l, ptr, len, unused); - } + memset (rule_buf_out, 0, sizeof (rule_buf_out)); - if (rule_len_out < 0) - { - continue; - } + const int rule_len_out = _old_apply_rule (user_options->rule_buf_l, user_options_extra->rule_len_l, ptr, len, rule_buf_out); - if (rule_len_out > PW_MAX) - { - continue; - } - } - else - { - if (len > PW_MAX) - { - continue; - } + if (rule_len_out < 0) continue; } + if (len >= PW_MAX) continue; + *out_buf = ptr; *out_len = len; @@ -387,6 +368,8 @@ int count_words (hashcat_ctx_t *hashcat_ctx, FILE *fd, const char *dictfile, u64 wl_data->func (ptr, wl_data->cnt - i, &len, &off); + i += off; + // do the on-the-fly encoding if (wl_data->iconv_enabled == true) @@ -398,59 +381,43 @@ int count_words (hashcat_ctx_t *hashcat_ctx, FILE *fd, const char *dictfile, u64 const size_t iconv_rc = iconv (wl_data->iconv_ctx, &ptr, &ptr_len, &iconv_ptr, &iconv_sz); - if (iconv_rc == (size_t) -1) - { - len = PW_MAX1; - } - else - { - ptr = wl_data->iconv_tmp; - len = HCBUFSIZ_TINY - iconv_sz; - } + if (iconv_rc == (size_t) -1) continue; + + ptr = wl_data->iconv_tmp; + len = HCBUFSIZ_TINY - iconv_sz; } if (run_rule_engine (user_options_extra->rule_len_l, user_options->rule_buf_l)) { - int rule_len_out = -1; - - if (len < BLOCK_SIZE) - { - char unused[BLOCK_SIZE] = { 0 }; - - rule_len_out = _old_apply_rule (user_options->rule_buf_l, user_options_extra->rule_len_l, ptr, len, unused); - } - - if (rule_len_out < 0) - { - len = PW_MAX1; - } - else - { - len = rule_len_out; - } - } + if (len >= BLOCK_SIZE) continue; - if (len < PW_MAX1) - { - if (user_options_extra->attack_kern == ATTACK_KERN_STRAIGHT) - { - if (overflow_check_u64_add (cnt, straight_ctx->kernel_rules_cnt) == false) return -1; + char rule_buf_out[BLOCK_SIZE]; - cnt += straight_ctx->kernel_rules_cnt; - } - else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) - { - if (overflow_check_u64_add (cnt, combinator_ctx->combs_cnt) == false) return -1; + memset (rule_buf_out, 0, sizeof (rule_buf_out)); - cnt += combinator_ctx->combs_cnt; - } + const int rule_len_out = _old_apply_rule (user_options->rule_buf_l, user_options_extra->rule_len_l, ptr, len, rule_buf_out); - d.cnt++; + if (rule_len_out < 0) continue; } - i += off; - cnt2++; + + if (len >= PW_MAX) continue; + + d.cnt++; + + if (user_options_extra->attack_kern == ATTACK_KERN_STRAIGHT) + { + if (overflow_check_u64_add (cnt, straight_ctx->kernel_rules_cnt) == false) return -1; + + cnt += straight_ctx->kernel_rules_cnt; + } + else if (user_options_extra->attack_kern == ATTACK_KERN_COMBI) + { + if (overflow_check_u64_add (cnt, combinator_ctx->combs_cnt) == false) return -1; + + cnt += combinator_ctx->combs_cnt; + } } time (&now);