arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-22 22:58:30 +00:00
Code Issues Releases Wiki Activity

Multiple changes:

Browse Source 
* 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
This commit is contained in:
jsteube 2017-06-29 12:19:05 +02:00
parent 6fb79b726c
commit f7a8e7c54b
17 changed files with 529 additions and 267 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

42
OpenCL/inc_scalar.cl Normal file
View File

@ -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); \
} \
} \
} \
}

98
OpenCL/m00000_a1.cl
View File

@ -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);
}
}

3
docs/changes.txt
View File

@ -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

8
extra/tab_completion/hashcat.sh
View File

@ -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
;;

8
include/common.h
View File

@ -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)

38
include/types.h
View File

@ -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;

17
src/autotune.c
View File

@ -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_1, kernel_power_try, true, 0);
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);
}
}
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));

37
src/combinator.c
View File

@ -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;

38
src/dispatch.c
View File

@ -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))
{
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;

159
src/interface.c
View File

@ -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 / 2;
hashconfig->pw_max = PW_MAX;
}
else
{
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;

151
src/opencl.c
View File

@ -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)
{
if (combinator_ctx->combs_mode == COMBINATOR_MODE_BASE_RIGHT)
}
else
{
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
}
}
}
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->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;
u8 *ptr = (u8 *) device_param->pws_buf[i].i;
ptr[pw_len] = 0x80;
}
}
}
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; }

25
src/outfile.c
View File

@ -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)
{

39
src/selftest.c
View File

@ -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;

4
src/straight.c
View File

@ -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)

2
src/usage.c
View File

@ -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",

14
src/user_options.c
View File

@ -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);

113
src/wordlist.c
View File

@ -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) 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)
{
len = PW_MAX1;
}
else
{
len = rule_len_out;
}
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 < 0) 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;
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;
}
d.cnt++;
}
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);