#define MAX_DICTSTAT 10000
-#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 136
+#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 143
#define global_free(attr) \
{ \
133,
13500,
11600,
+ 13600,
12500,
13000,
13200,
6221,
6231,
6241,
+ 13711,
+ 13721,
+ 13731,
+ 13741,
+ 13751,
+ 13761,
8800,
12900,
12200,
" --outfile-check-dir=FOLDER Specify the outfile directory which should be monitored, default is $session.outfiles",
" --logfile-disable Disable the logfile",
" --truecrypt-keyfiles=FILE Keyfiles used, separate with comma",
+ " --veracrypt-keyfiles=FILE Keyfiles used, separate with comma",
+ " --veracrypt-pim=NUM VeraCrypt personal iterations multiplier",
"",
"* Resources:",
"",
" 13000 = RAR5",
" 13200 = AxCrypt",
" 13300 = AxCrypt in memory SHA1",
+ " 13600 = WinZip",
"",
"[[ Full-Disk encryptions (FDE) ]]",
"",
- " 62XY = TrueCrypt 5.0+",
+ " 62XY = TrueCrypt",
" X = 1 = PBKDF2-HMAC-RipeMD160",
" X = 2 = PBKDF2-HMAC-SHA512",
" X = 3 = PBKDF2-HMAC-Whirlpool",
" 8800 = Android FDE < v4.3",
" 12900 = Android FDE (Samsung DEK)",
" 12200 = eCryptfs",
+ " 137XY = VeraCrypt",
+ " X = 1 = PBKDF2-HMAC-RipeMD160",
+ " X = 2 = PBKDF2-HMAC-SHA512",
+ " X = 3 = PBKDF2-HMAC-Whirlpool",
+ " X = 4 = PBKDF2-HMAC-RipeMD160 + boot-mode",
+ " X = 5 = PBKDF2-HMAC-SHA256",
+ " X = 6 = PBKDF2-HMAC-SHA256 + boot-mode",
+ " Y = 1 = XTS 512 bit (Ciphers: AES or Serpent or Twofish)",
+ " Y = 2 = XTS 1024 bit (Ciphers: AES or Serpent or Twofish or AES-Twofish or Serpent-AES or Twofish-Serpent)",
+ " Y = 3 = XTS 1536 bit (Ciphers: All)",
"",
"[[ Documents ]]",
"",
{
log_info ("Hash.Target....: File (%s)", data.hashfile);
}
+ else if ((data.hash_mode >= 13700) && (data.hash_mode <= 13799))
+ {
+ log_info ("Hash.Target....: File (%s)", data.hashfile);
+ }
else
{
char out_buf[HCBUFSIZ] = { 0 };
speed_ms[device_id] /= SPEED_CACHE;
}
- float hashes_all_ms = 0;
+ double hashes_all_ms = 0;
- float hashes_dev_ms[DEVICES_MAX] = { 0 };
+ double hashes_dev_ms[DEVICES_MAX] = { 0 };
for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
{
if (speed_ms[device_id])
{
- hashes_dev_ms[device_id] = speed_cnt[device_id] / speed_ms[device_id];
+ hashes_dev_ms[device_id] = (double) speed_cnt[device_id] / speed_ms[device_id];
hashes_all_ms += hashes_dev_ms[device_id];
}
#endif // HAVE_HWMON
}
+static void status_benchmark_automat ()
+{
+ u64 speed_cnt[DEVICES_MAX] = { 0 };
+ double speed_ms[DEVICES_MAX] = { 0 };
+
+ for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
+ {
+ hc_device_param_t *device_param = &data.devices_param[device_id];
+
+ if (device_param->skipped) continue;
+
+ speed_cnt[device_id] = device_param->speed_cnt[0];
+ speed_ms[device_id] = device_param->speed_ms[0];
+ }
+
+ double hashes_dev_ms[DEVICES_MAX] = { 0 };
+
+ for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
+ {
+ hc_device_param_t *device_param = &data.devices_param[device_id];
+
+ if (device_param->skipped) continue;
+
+ hashes_dev_ms[device_id] = 0;
+
+ if (speed_ms[device_id])
+ {
+ hashes_dev_ms[device_id] = (double) speed_cnt[device_id] / speed_ms[device_id];
+ }
+ }
+
+ for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
+ {
+ hc_device_param_t *device_param = &data.devices_param[device_id];
+
+ if (device_param->skipped) continue;
+
+ log_info ("%u:%u:%llu", device_id + 1, data.hash_mode, (unsigned long long int) (hashes_dev_ms[device_id] * 1000));
+ }
+}
+
static void status_benchmark ()
{
if (data.devices_status == STATUS_INIT) return;
if (data.devices_status == STATUS_STARTING) return;
+ if (data.devices_status == STATUS_BYPASS) return;
- if (data.words_cnt == 0) return;
+ if (data.status_automat == 1)
+ {
+ status_benchmark_automat ();
+
+ return;
+ }
u64 speed_cnt[DEVICES_MAX] = { 0 };
double speed_ms[DEVICES_MAX] = { 0 };
speed_ms[device_id] = device_param->speed_ms[0];
}
- float hashes_all_ms = 0;
+ double hashes_all_ms = 0;
- float hashes_dev_ms[DEVICES_MAX] = { 0 };
+ double hashes_dev_ms[DEVICES_MAX] = { 0 };
for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
{
if (speed_ms[device_id])
{
- hashes_dev_ms[device_id] = speed_cnt[device_id] / speed_ms[device_id];
+ hashes_dev_ms[device_id] = (double) speed_cnt[device_id] / speed_ms[device_id];
hashes_all_ms += hashes_dev_ms[device_id];
}
hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, gidd * sizeof (pw_t), sizeof (pw_t), pw, 0, NULL, NULL);
}
-static void check_hash (hc_device_param_t *device_param, const uint salt_pos, const uint digest_pos)
+static void check_hash (hc_device_param_t *device_param, plain_t *plain)
{
char *outfile = data.outfile;
uint quiet = data.quiet;
char out_buf[HCBUFSIZ] = { 0 };
- ascii_digest (out_buf, salt_pos, digest_pos);
+ const u32 salt_pos = plain->salt_pos;
+ const u32 digest_pos = plain->digest_pos; // relative
+ const u32 gidvid = plain->gidvid;
+ const u32 il_pos = plain->il_pos;
- uint idx = data.salts_buf[salt_pos].digests_offset + digest_pos;
+ ascii_digest (out_buf, salt_pos, digest_pos);
// plain
- plain_t plain;
-
- hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, idx * sizeof (plain_t), sizeof (plain_t), &plain, 0, NULL, NULL);
-
- uint gidvid = plain.gidvid;
- uint il_pos = plain.il_pos;
-
u64 crackpos = device_param->words_off;
uint plain_buf[16] = { 0 };
u8 *plain_ptr = (u8 *) plain_buf;
+
unsigned int plain_len = 0;
if (data.attack_mode == ATTACK_MODE_STRAIGHT)
{
- u64 gidd = gidvid;
- u64 gidm = 0;
-
pw_t pw;
- gidd_to_pw_t (device_param, gidd, &pw);
+ gidd_to_pw_t (device_param, gidvid, &pw);
- for (int i = 0, j = gidm; i < 16; i++, j++)
+ for (int i = 0; i < 16; i++)
{
- plain_buf[i] = pw.i[j];
+ plain_buf[i] = pw.i[i];
}
plain_len = pw.pw_len;
}
else if (data.attack_mode == ATTACK_MODE_COMBI)
{
- u64 gidd = gidvid;
- u64 gidm = 0;
-
pw_t pw;
- gidd_to_pw_t (device_param, gidd, &pw);
+ gidd_to_pw_t (device_param, gidvid, &pw);
- for (int i = 0, j = gidm; i < 16; i++, j++)
+ for (int i = 0; i < 16; i++)
{
- plain_buf[i] = pw.i[j];
+ plain_buf[i] = pw.i[i];
}
plain_len = pw.pw_len;
}
else if (data.attack_mode == ATTACK_MODE_HYBRID1)
{
- u64 gidd = gidvid;
- u64 gidm = 0;
-
pw_t pw;
- gidd_to_pw_t (device_param, gidd, &pw);
+ gidd_to_pw_t (device_param, gidvid, &pw);
- for (int i = 0, j = gidm; i < 16; i++, j++)
+ for (int i = 0; i < 16; i++)
{
- plain_buf[i] = pw.i[j];
+ plain_buf[i] = pw.i[i];
}
plain_len = pw.pw_len;
}
else if (data.attack_mode == ATTACK_MODE_HYBRID2)
{
- u64 gidd = gidvid;
- u64 gidm = 0;
-
pw_t pw;
- gidd_to_pw_t (device_param, gidd, &pw);
+ gidd_to_pw_t (device_param, gidvid, &pw);
- for (int i = 0, j = gidm; i < 16; i++, j++)
+ for (int i = 0; i < 16; i++)
{
- plain_buf[i] = pw.i[j];
+ plain_buf[i] = pw.i[i];
}
plain_len = pw.pw_len;
if ((quiet == 0) && (debug_file == NULL))
{
fprintf (stdout, "%s", PROMPT);
+
fflush (stdout);
}
}
{
salt_t *salt_buf = &data.salts_buf[salt_pos];
- int found = 0;
+ u32 num_cracked;
- hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL);
- for (uint i = 0; i < device_param->kernel_threads; i++) if (device_param->result[i] == 1) found = 1;
-
- if (found == 1)
+ if (num_cracked)
{
// display hack (for weak hashes etc, it could be that there is still something to clear on the current line)
log_info_nn ("");
- hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
+ plain_t *cracked = (plain_t *) mycalloc (num_cracked, sizeof (plain_t));
+
+ hc_clEnqueueReadBuffer (data.ocl, device_param->command_queue, device_param->d_plain_bufs, CL_TRUE, 0, num_cracked * sizeof (plain_t), cracked, 0, NULL, NULL);
uint cpt_cracked = 0;
- for (uint digest_pos = 0; digest_pos < salt_buf->digests_cnt; digest_pos++)
+ for (uint i = 0; i < num_cracked; i++)
{
- uint idx = salt_buf->digests_offset + digest_pos;
+ const uint hash_pos = cracked[i].hash_pos;
- if (data.digests_shown_tmp[idx] == 0) continue;
-
- if (data.digests_shown[idx] == 1) continue;
+ if (data.digests_shown[hash_pos] == 1) continue;
if ((data.opts_type & OPTS_TYPE_PT_NEVERCRACK) == 0)
{
- data.digests_shown[idx] = 1;
+ data.digests_shown[hash_pos] = 1;
data.digests_done++;
if (data.salts_done == data.salts_cnt) data.devices_status = STATUS_CRACKED;
- check_hash (device_param, salt_pos, digest_pos);
+ check_hash (device_param, &cracked[i]);
}
+ myfree (cracked);
+
if (cpt_cracked > 0)
{
data.cpt_buf[data.cpt_pos].timestamp = time (NULL);
hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, salt_buf->digests_offset * sizeof (uint), salt_buf->digests_cnt * sizeof (uint), &data.digests_shown_tmp[salt_buf->digests_offset], 0, NULL, NULL);
}
- memset (device_param->result, 0, device_param->size_results);
+ num_cracked = 0;
- hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, device_param->size_results, device_param->result, 0, NULL, NULL);
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_result, CL_TRUE, 0, sizeof (u32), &num_cracked, 0, NULL, NULL);
}
}
{
clear_prompt ();
- log_info ("");
+ //log_info ("");
log_info ("INFO: approaching final keyspace, workload adjusted");
-
log_info ("");
fprintf (stdout, "%s", PROMPT);
hc_clSetKernelArg (data.ocl, kernel, 30, sizeof (cl_uint), device_param->kernel_params[30]);
hc_clSetKernelArg (data.ocl, kernel, 31, sizeof (cl_uint), device_param->kernel_params[31]);
- hc_timer_t timer;
-
- hc_timer_set (&timer);
+ cl_event event;
if ((data.opts_type & OPTS_TYPE_PT_BITSLICE) && (data.attack_mode == ATTACK_MODE_BF))
{
const size_t global_work_size[3] = { num_elements, 32, 1 };
const size_t local_work_size[3] = { kernel_threads / 32, 32, 1 };
- hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 2, NULL, global_work_size, local_work_size, 0, NULL, &event);
}
else
{
- size_t workgroup_size = 0;
-
- hc_clGetKernelWorkGroupInfo (data.ocl, kernel, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &workgroup_size, NULL);
-
if (kern_run == KERN_RUN_2)
{
if (data.opti_type & OPTI_TYPE_SLOW_HASH_SIMD)
}
}
- if (kernel_threads > workgroup_size) kernel_threads = workgroup_size;
-
while (num_elements % kernel_threads) num_elements++;
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
- hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, NULL);
+ hc_clEnqueueNDRangeKernel (data.ocl, device_param->command_queue, kernel, 1, NULL, global_work_size, local_work_size, 0, NULL, &event);
}
hc_clFlush (data.ocl, device_param->command_queue);
- hc_clFinish (data.ocl, device_param->command_queue);
+ hc_clWaitForEvents (data.ocl, 1, &event);
if (event_update)
{
- double exec_time;
+ cl_ulong time_start;
+ cl_ulong time_end;
- hc_timer_get (timer, exec_time);
+ hc_clGetEventProfilingInfo (data.ocl, event, CL_PROFILING_COMMAND_START, sizeof (time_start), &time_start, NULL);
+ hc_clGetEventProfilingInfo (data.ocl, event, CL_PROFILING_COMMAND_END, sizeof (time_end), &time_end, NULL);
+
+ const double exec_time = (double) (time_end - time_start) / 1000000.0;
uint exec_pos = device_param->exec_pos;
device_param->exec_pos = exec_pos;
}
+
+ hc_clReleaseEvent (data.ocl, event);
+
+ hc_clFinish (data.ocl, device_param->command_queue);
}
static void run_kernel_mp (const uint kern_run, hc_device_param_t *device_param, const uint num)
break;
}
- size_t workgroup_size = 0;
-
- hc_clGetKernelWorkGroupInfo (data.ocl, kernel, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(size_t), &workgroup_size, NULL);
-
- if (kernel_threads > workgroup_size) kernel_threads = workgroup_size;
-
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
cl_kernel kernel = device_param->kernel_tm;
- size_t workgroup_size = 0;
-
- hc_clGetKernelWorkGroupInfo (data.ocl, kernel, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &workgroup_size, NULL);
-
- if (kernel_threads > workgroup_size) kernel_threads = workgroup_size;
-
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
hc_clSetKernelArg (data.ocl, kernel, 5, sizeof (cl_uint), device_param->kernel_params_amp[5]);
hc_clSetKernelArg (data.ocl, kernel, 6, sizeof (cl_uint), device_param->kernel_params_amp[6]);
- size_t workgroup_size = 0;
-
- hc_clGetKernelWorkGroupInfo (data.ocl, kernel, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &workgroup_size, NULL);
-
- if (kernel_threads > workgroup_size) kernel_threads = workgroup_size;
-
const size_t global_work_size[3] = { num_elements, 1, 1 };
const size_t local_work_size[3] = { kernel_threads, 1, 1 };
}
}
- // sometimes we're in a bad situation that the algorithm is so slow that we can not
- // create enough kernel_accel to do both, keep the gpu busy and stay below target_ms.
- // however, we need to have a minimum kernel_accel and kernel_loops of 32.
- // luckily, at this level of workload, it became a linear function
+ // balancing the workload turns out to be very efficient
- if (kernel_accel < 32 || kernel_loops < 32)
+ if (kernel_loops_min != kernel_loops_max)
{
- const u32 kernel_power = kernel_accel * kernel_loops;
-
- // find sqrt
+ const u32 kernel_power_balance = kernel_accel * kernel_loops;
u32 sqrtv;
for (sqrtv = 1; sqrtv < 0x100000; sqrtv++)
{
- if ((sqrtv * sqrtv) >= kernel_power) break;
+ if ((sqrtv * sqrtv) >= kernel_power_balance) break;
}
const u32 kernel_accel_try = sqrtv;
device_param->pws_cnt = 0;
+ /*
+ still required?
if (attack_kern == ATTACK_KERN_STRAIGHT)
{
run_kernel_bzero (device_param, device_param->d_rules_c, device_param->size_rules_c);
{
run_kernel_bzero (device_param, device_param->d_combs_c, device_param->size_combs);
}
+ */
}
}
device_param->pws_cnt = 0;
+ /*
+ still required?
run_kernel_bzero (device_param, device_param->d_bfs_c, device_param->size_bfs);
+ */
}
if (data.devices_status == STATUS_STOP_AT_CHECKPOINT) check_checkpoint ();
device_param->pws_cnt = 0;
+ /*
+ still required?
if (attack_kern == ATTACK_KERN_STRAIGHT)
{
run_kernel_bzero (device_param, device_param->d_rules_c, device_param->size_rules_c);
{
run_kernel_bzero (device_param, device_param->d_combs_c, device_param->size_combs);
}
+ */
}
if (data.devices_status == STATUS_STOP_AT_CHECKPOINT) check_checkpoint ();
for (uint i = 0; i < digests_cnt; i++)
{
+ if (data.digests_shown[i] == 1) continue; // can happen with potfile
+
uint *digest_ptr = (uint *) digests_buf_ptr;
digests_buf_ptr += dgst_size;
if (getenv ("POCL_KERNEL_CACHE") == NULL)
putenv ((char *) "POCL_KERNEL_CACHE=0");
+ umask (077);
+
/**
* Real init
*/
char *opencl_device_types = NULL;
uint opencl_vector_width = OPENCL_VECTOR_WIDTH;
char *truecrypt_keyfiles = NULL;
+ char *veracrypt_keyfiles = NULL;
+ uint veracrypt_pim = 0;
uint workload_profile = WORKLOAD_PROFILE;
uint kernel_accel = KERNEL_ACCEL;
uint kernel_loops = KERNEL_LOOPS;
#define IDX_POWERTUNE_ENABLE 0xff41
#define IDX_LOGFILE_DISABLE 0xff51
#define IDX_TRUECRYPT_KEYFILES 0xff52
+ #define IDX_VERACRYPT_KEYFILES 0xff53
+ #define IDX_VERACRYPT_PIM 0xff54
#define IDX_SCRYPT_TMTO 0xff61
#define IDX_SEGMENT_SIZE 'c'
#define IDX_SEPARATOR 'p'
#endif // HAVE_HWMON
{"logfile-disable", no_argument, 0, IDX_LOGFILE_DISABLE},
{"truecrypt-keyfiles", required_argument, 0, IDX_TRUECRYPT_KEYFILES},
+ {"veracrypt-keyfiles", required_argument, 0, IDX_VERACRYPT_KEYFILES},
+ {"veracrypt-pim", required_argument, 0, IDX_VERACRYPT_PIM},
{"segment-size", required_argument, 0, IDX_SEGMENT_SIZE},
{"scrypt-tmto", required_argument, 0, IDX_SCRYPT_TMTO},
// deprecated
if (version)
{
- log_info ("%s (%s)", VERSION_TAG, VERSION_SUM);
+ log_info ("%s", VERSION_TAG);
return (0);
}
#endif // HAVE_HWMON
case IDX_LOGFILE_DISABLE: logfile_disable = 1; break;
case IDX_TRUECRYPT_KEYFILES: truecrypt_keyfiles = optarg; break;
+ case IDX_VERACRYPT_KEYFILES: veracrypt_keyfiles = optarg; break;
+ case IDX_VERACRYPT_PIM: veracrypt_pim = atoi (optarg); break;
case IDX_SEGMENT_SIZE: segment_size = atoi (optarg); break;
case IDX_SCRYPT_TMTO: scrypt_tmto = atoi (optarg); break;
case IDX_SEPARATOR: separator = optarg[0]; break;
{
if (benchmark == 1)
{
- log_info ("%s %s (%s) starting in benchmark-mode...", PROGNAME, VERSION_TAG, VERSION_SUM);
-
- log_info ("");
+ if (status_automat == 0)
+ {
+ log_info ("%s (%s) starting in benchmark-mode...", PROGNAME, VERSION_TAG);
+ log_info ("");
+ }
+ else
+ {
+ log_info ("# %s (%s) %s", PROGNAME, VERSION_TAG, ctime (&proc_start));
+ }
}
else if (restore == 1)
{
- log_info ("%s %s (%s) starting in restore-mode...", PROGNAME, VERSION_TAG, VERSION_SUM);
-
+ log_info ("%s (%s) starting in restore-mode...", PROGNAME, VERSION_TAG);
log_info ("");
}
else
{
- log_info ("%s %s (%s) starting...", PROGNAME, VERSION_TAG, VERSION_SUM);
-
+ log_info ("%s (%s) starting...", PROGNAME, VERSION_TAG);
log_info ("");
}
}
return (-1);
}
- if (hash_mode_chgd && hash_mode > 13500) // just added to remove compiler warnings for hash_mode_chgd
+ if (hash_mode_chgd && hash_mode > 13799) // just added to remove compiler warnings for hash_mode_chgd
{
log_error ("ERROR: Invalid hash-type specified");
if (username == 1)
{
- if ((hash_mode == 2500) || (hash_mode == 5200) || ((hash_mode >= 6200) && (hash_mode <= 6299)))
+ if ((hash_mode == 2500) || (hash_mode == 5200) || ((hash_mode >= 6200) && (hash_mode <= 6299)) || ((hash_mode >= 13700) && (hash_mode <= 13799)))
{
log_error ("ERROR: Mixing support for user names and hashes of type %s is not supported", strhashtype (hash_mode));
#endif
data.logfile_disable = logfile_disable;
data.truecrypt_keyfiles = truecrypt_keyfiles;
+ data.veracrypt_keyfiles = veracrypt_keyfiles;
+ data.veracrypt_pim = veracrypt_pim;
data.scrypt_tmto = scrypt_tmto;
data.workload_profile = workload_profile;
logfile_top_string (rule_buf_r);
logfile_top_string (session);
logfile_top_string (truecrypt_keyfiles);
+ logfile_top_string (veracrypt_keyfiles);
+ logfile_top_uint (veracrypt_pim);
/**
* Init OpenCL library loader
dgst_size = DGST_SIZE_4_4;
parse_func = dcc2_parse_hash;
sort_by_digest = sort_by_digest_4_4;
- opti_type = OPTI_TYPE_ZERO_BYTE;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
dgst_size = DGST_SIZE_4_4;
parse_func = wpa_parse_hash;
sort_by_digest = sort_by_digest_4_4;
- opti_type = OPTI_TYPE_ZERO_BYTE;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
parse_func = sha512osx_parse_hash;
sort_by_digest = sort_by_digest_8_16;
opti_type = OPTI_TYPE_ZERO_BYTE
- | OPTI_TYPE_USES_BITS_64;
+ | OPTI_TYPE_USES_BITS_64
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
parse_func = sha512grub_parse_hash;
sort_by_digest = sort_by_digest_8_16;
opti_type = OPTI_TYPE_ZERO_BYTE
- | OPTI_TYPE_USES_BITS_64;
+ | OPTI_TYPE_USES_BITS_64
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
dgst_size = DGST_SIZE_4_32;
parse_func = cisco8_parse_hash;
sort_by_digest = sort_by_digest_4_32;
- opti_type = OPTI_TYPE_ZERO_BYTE;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
dgst_size = DGST_SIZE_4_32;
parse_func = djangopbkdf2_parse_hash;
sort_by_digest = sort_by_digest_4_32;
- opti_type = OPTI_TYPE_ZERO_BYTE;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
dgst_size = DGST_SIZE_4_32;
parse_func = pbkdf2_sha256_parse_hash;
sort_by_digest = sort_by_digest_4_32;
- opti_type = OPTI_TYPE_ZERO_BYTE;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
dgst_size = DGST_SIZE_4_32;
parse_func = pbkdf2_md5_parse_hash;
sort_by_digest = sort_by_digest_4_32;
- opti_type = OPTI_TYPE_ZERO_BYTE;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
dgst_size = DGST_SIZE_4_32;
parse_func = pbkdf2_sha1_parse_hash;
sort_by_digest = sort_by_digest_4_32;
- opti_type = OPTI_TYPE_ZERO_BYTE;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
parse_func = pbkdf2_sha512_parse_hash;
sort_by_digest = sort_by_digest_8_16;
opti_type = OPTI_TYPE_ZERO_BYTE
- | OPTI_TYPE_USES_BITS_64;
+ | OPTI_TYPE_USES_BITS_64
+ | OPTI_TYPE_SLOW_HASH_SIMD;
dgst_pos0 = 0;
dgst_pos1 = 1;
dgst_pos2 = 2;
dgst_pos3 = 1;
break;
- default: usage_mini_print (PROGNAME); return (-1);
- }
-
- /**
- * parser
- */
+ case 13600: hash_type = HASH_TYPE_PBKDF2_SHA1;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_ZIP2;
+ dgst_size = DGST_SIZE_4_4;
+ parse_func = zip2_parse_hash;
+ sort_by_digest = sort_by_digest_4_4;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
- data.parse_func = parse_func;
+ case 13711: hash_type = HASH_TYPE_RIPEMD160;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCRIPEMD160_XTS512;
+ dgst_size = DGST_SIZE_4_5;
+ parse_func = veracrypt_parse_hash_655331;
+ sort_by_digest = sort_by_digest_4_5;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
- /**
- * misc stuff
- */
+ case 13712: hash_type = HASH_TYPE_RIPEMD160;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCRIPEMD160_XTS1024;
+ dgst_size = DGST_SIZE_4_5;
+ parse_func = veracrypt_parse_hash_655331;
+ sort_by_digest = sort_by_digest_4_5;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
- if (hex_salt)
- {
- if (salt_type == SALT_TYPE_INTERN)
- {
- opts_type |= OPTS_TYPE_ST_HEX;
- }
- else
- {
- log_error ("ERROR: Parameter hex-salt not valid for hash-type %u", hash_mode);
+ case 13713: hash_type = HASH_TYPE_RIPEMD160;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCRIPEMD160_XTS1536;
+ dgst_size = DGST_SIZE_4_5;
+ parse_func = veracrypt_parse_hash_655331;
+ sort_by_digest = sort_by_digest_4_5;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
- return (-1);
- }
- }
+ case 13721: hash_type = HASH_TYPE_SHA512;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_TCSHA512_XTS512;
+ dgst_size = DGST_SIZE_8_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_8_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_USES_BITS_64;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
- uint isSalted = ((salt_type == SALT_TYPE_INTERN)
- | (salt_type == SALT_TYPE_EXTERN)
- | (salt_type == SALT_TYPE_EMBEDDED)
- | (salt_type == SALT_TYPE_VIRTUAL));
+ case 13722: hash_type = HASH_TYPE_SHA512;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_TCSHA512_XTS1024;
+ dgst_size = DGST_SIZE_8_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_8_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_USES_BITS_64;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
- sort_by_digest = sort_by_digest_p0p1; // overruled by 64 bit digest
+ case 13723: hash_type = HASH_TYPE_SHA512;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_TCSHA512_XTS1536;
+ dgst_size = DGST_SIZE_8_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_8_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_USES_BITS_64;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
- data.hash_type = hash_type;
- data.attack_mode = attack_mode;
- data.attack_kern = attack_kern;
+ case 13731: hash_type = HASH_TYPE_WHIRLPOOL;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCWHIRLPOOL_XTS512;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13732: hash_type = HASH_TYPE_WHIRLPOOL;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCWHIRLPOOL_XTS1024;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13733: hash_type = HASH_TYPE_WHIRLPOOL;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCWHIRLPOOL_XTS1536;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13741: hash_type = HASH_TYPE_RIPEMD160;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCRIPEMD160_XTS512;
+ dgst_size = DGST_SIZE_4_5;
+ parse_func = veracrypt_parse_hash_327661;
+ sort_by_digest = sort_by_digest_4_5;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13742: hash_type = HASH_TYPE_RIPEMD160;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCRIPEMD160_XTS1024;
+ dgst_size = DGST_SIZE_4_5;
+ parse_func = veracrypt_parse_hash_327661;
+ sort_by_digest = sort_by_digest_4_5;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13743: hash_type = HASH_TYPE_RIPEMD160;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_TCRIPEMD160_XTS1536;
+ dgst_size = DGST_SIZE_4_5;
+ parse_func = veracrypt_parse_hash_327661;
+ sort_by_digest = sort_by_digest_4_5;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13751: hash_type = HASH_TYPE_SHA256;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_VCSHA256_XTS512;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13752: hash_type = HASH_TYPE_SHA256;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_VCSHA256_XTS1024;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13753: hash_type = HASH_TYPE_SHA256;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_VCSHA256_XTS1536;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_500000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13761: hash_type = HASH_TYPE_SHA256;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_VCSHA256_XTS512;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_200000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13762: hash_type = HASH_TYPE_SHA256;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_VCSHA256_XTS1024;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_200000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ case 13763: hash_type = HASH_TYPE_SHA256;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE; // should be OPTS_TYPE_PT_GENERATE_BE
+ kern_type = KERN_TYPE_VCSHA256_XTS1536;
+ dgst_size = DGST_SIZE_4_8;
+ parse_func = veracrypt_parse_hash_200000;
+ sort_by_digest = sort_by_digest_4_8;
+ opti_type = OPTI_TYPE_ZERO_BYTE;
+ dgst_pos0 = 0;
+ dgst_pos1 = 1;
+ dgst_pos2 = 2;
+ dgst_pos3 = 3;
+ break;
+
+ default: usage_mini_print (PROGNAME); return (-1);
+ }
+
+ /**
+ * parser
+ */
+
+ data.parse_func = parse_func;
+
+ /**
+ * misc stuff
+ */
+
+ if (hex_salt)
+ {
+ if (salt_type == SALT_TYPE_INTERN)
+ {
+ opts_type |= OPTS_TYPE_ST_HEX;
+ }
+ else
+ {
+ log_error ("ERROR: Parameter hex-salt not valid for hash-type %u", hash_mode);
+
+ return (-1);
+ }
+ }
+
+ uint isSalted = ((salt_type == SALT_TYPE_INTERN)
+ | (salt_type == SALT_TYPE_EXTERN)
+ | (salt_type == SALT_TYPE_EMBEDDED)
+ | (salt_type == SALT_TYPE_VIRTUAL));
+
+ sort_by_digest = sort_by_digest_p0p1; // overruled by 64 bit digest
+
+ data.hash_type = hash_type;
+ data.attack_mode = attack_mode;
+ data.attack_kern = attack_kern;
data.attack_exec = attack_exec;
data.kern_type = kern_type;
data.opts_type = opts_type;
case 13100: esalt_size = sizeof (krb5tgs_t); break;
case 13400: esalt_size = sizeof (keepass_t); break;
case 13500: esalt_size = sizeof (pstoken_t); break;
+ case 13600: esalt_size = sizeof (zip2_t); break;
+ case 13711: esalt_size = sizeof (tc_t); break;
+ case 13712: esalt_size = sizeof (tc_t); break;
+ case 13713: esalt_size = sizeof (tc_t); break;
+ case 13721: esalt_size = sizeof (tc_t); break;
+ case 13722: esalt_size = sizeof (tc_t); break;
+ case 13723: esalt_size = sizeof (tc_t); break;
+ case 13731: esalt_size = sizeof (tc_t); break;
+ case 13732: esalt_size = sizeof (tc_t); break;
+ case 13733: esalt_size = sizeof (tc_t); break;
+ case 13741: esalt_size = sizeof (tc_t); break;
+ case 13742: esalt_size = sizeof (tc_t); break;
+ case 13743: esalt_size = sizeof (tc_t); break;
+ case 13751: esalt_size = sizeof (tc_t); break;
+ case 13752: esalt_size = sizeof (tc_t); break;
+ case 13753: esalt_size = sizeof (tc_t); break;
+ case 13761: esalt_size = sizeof (tc_t); break;
+ case 13762: esalt_size = sizeof (tc_t); break;
+ case 13763: esalt_size = sizeof (tc_t); break;
}
data.esalt_size = esalt_size;
if ((hash_mode == 2500) ||
(hash_mode == 5200) ||
- ((hash_mode >= 6200) && (hash_mode <= 6299)) ||
+ ((hash_mode >= 6200) && (hash_mode <= 6299)) ||
+ ((hash_mode >= 13700) && (hash_mode <= 13799)) ||
(hash_mode == 9000))
{
hashlist_mode = HL_MODE_ARG;
((seven_zip_t *) hashes_buf[0].esalt)->data_len = 112;
((seven_zip_t *) hashes_buf[0].esalt)->unpack_size = 112;
break;
- case 13400: ((keepass_t *) hashes_buf[0].esalt)->version = 2;
+ case 13400: ((keepass_t *) hashes_buf[0].esalt)->version = 2;
+ break;
+ case 13500: ((pstoken_t *) hashes_buf[0].esalt)->salt_len = 113;
break;
- case 13500: ((pstoken_t *) hashes_buf[0].esalt)->salt_len = 113;
+ case 13600: ((zip2_t *) hashes_buf[0].esalt)->salt_len = 16;
+ ((zip2_t *) hashes_buf[0].esalt)->data_len = 32;
+ ((zip2_t *) hashes_buf[0].esalt)->mode = 3;
break;
}
}
break;
case 9000: data.hashfile = mystrdup ("hashcat.psafe2");
break;
+ case 13711: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13712: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13713: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13721: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13722: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13723: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13731: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13732: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13733: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13741: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13742: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13743: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13751: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13752: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13753: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13761: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13762: data.hashfile = mystrdup ("hashcat.vc");
+ break;
+ case 13763: data.hashfile = mystrdup ("hashcat.vc");
+ break;
}
// set default iterations
break;
case 13400: hashes_buf[0].salt->salt_iter = ROUNDS_KEEPASS;
break;
+ case 13600: hashes_buf[0].salt->salt_iter = ROUNDS_ZIP2;
+ break;
+ case 13711: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_655331;
+ break;
+ case 13712: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_655331;
+ break;
+ case 13713: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_655331;
+ break;
+ case 13721: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13722: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13723: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13731: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13732: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13733: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13741: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_327661;
+ break;
+ case 13742: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_327661;
+ break;
+ case 13743: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_327661;
+ break;
+ case 13751: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13752: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13753: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_500000;
+ break;
+ case 13761: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_200000;
+ break;
+ case 13762: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_200000;
+ break;
+ case 13763: hashes_buf[0].salt->salt_iter = ROUNDS_VERACRYPT_200000;
+ break;
}
hashes_cnt = 1;
// no solution for these special hash types (for instane because they use hashfile in output etc)
if ((hash_mode != 5200) &&
- !((hash_mode >= 6200) && (hash_mode <= 6299)) &&
+ !((hash_mode >= 6200) && (hash_mode <= 6299)) &&
+ !((hash_mode >= 13700) && (hash_mode <= 13799)) &&
(hash_mode != 9000))
{
FILE *fp = fopen (potfile, "rb");
case 6241: salts_buf->truecrypt_mdlen = 1 * 512; break;
case 6242: salts_buf->truecrypt_mdlen = 2 * 512; break;
case 6243: salts_buf->truecrypt_mdlen = 3 * 512; break;
+ case 13711: salts_buf->truecrypt_mdlen = 1 * 512; break;
+ case 13712: salts_buf->truecrypt_mdlen = 2 * 512; break;
+ case 13713: salts_buf->truecrypt_mdlen = 3 * 512; break;
+ case 13721: salts_buf->truecrypt_mdlen = 1 * 512; break;
+ case 13722: salts_buf->truecrypt_mdlen = 2 * 512; break;
+ case 13723: salts_buf->truecrypt_mdlen = 3 * 512; break;
+ case 13731: salts_buf->truecrypt_mdlen = 1 * 512; break;
+ case 13732: salts_buf->truecrypt_mdlen = 2 * 512; break;
+ case 13733: salts_buf->truecrypt_mdlen = 3 * 512; break;
+ case 13741: salts_buf->truecrypt_mdlen = 1 * 512; break;
+ case 13742: salts_buf->truecrypt_mdlen = 2 * 512; break;
+ case 13743: salts_buf->truecrypt_mdlen = 3 * 512; break;
+ case 13751: salts_buf->truecrypt_mdlen = 1 * 512; break;
+ case 13752: salts_buf->truecrypt_mdlen = 2 * 512; break;
+ case 13753: salts_buf->truecrypt_mdlen = 3 * 512; break;
+ case 13761: salts_buf->truecrypt_mdlen = 1 * 512; break;
+ case 13762: salts_buf->truecrypt_mdlen = 2 * 512; break;
+ case 13763: salts_buf->truecrypt_mdlen = 3 * 512; break;
}
if (truecrypt_keyfiles)
free (keyfiles);
}
+ if (veracrypt_keyfiles)
+ {
+ uint *keyfile_buf = ((tc_t *) esalts_buf)->keyfile_buf;
+
+ char *keyfiles = strdup (veracrypt_keyfiles);
+
+ char *keyfile = strtok (keyfiles, ",");
+
+ do
+ {
+ truecrypt_crc32 (keyfile, (u8 *) keyfile_buf);
+
+ } while ((keyfile = strtok (NULL, ",")) != NULL);
+
+ free (keyfiles);
+ }
+
data.digests_cnt = digests_cnt;
data.digests_done = digests_done;
data.digests_buf = digests_buf;
}
else if (strcmp (platform_vendor, CL_VENDOR_APPLE) == 0)
{
- vendor_id = VENDOR_ID_GENERIC;
+ vendor_id = VENDOR_ID_APPLE;
}
else if (strcmp (platform_vendor, CL_VENDOR_INTEL_BEIGNET) == 0)
{
- vendor_id = VENDOR_ID_GENERIC;
+ vendor_id = VENDOR_ID_INTEL_BEIGNET;
}
else if (strcmp (platform_vendor, CL_VENDOR_INTEL_SDK) == 0)
{
- vendor_id = VENDOR_ID_GENERIC;
+ vendor_id = VENDOR_ID_INTEL_SDK;
}
else if (strcmp (platform_vendor, CL_VENDOR_MESA) == 0)
{
- vendor_id = VENDOR_ID_GENERIC;
+ vendor_id = VENDOR_ID_MESA;
}
else if (strcmp (platform_vendor, CL_VENDOR_NV) == 0)
{
}
else if (strcmp (platform_vendor, CL_VENDOR_POCL) == 0)
{
- vendor_id = VENDOR_ID_GENERIC;
+ vendor_id = VENDOR_ID_POCL;
}
else
{
if ((benchmark == 1 || quiet == 0) && (algorithm_pos == 0))
{
- if (device_param->skipped == 0)
+ if (status_automat == 0)
{
- log_info ("Device #%u: %s, %lu/%lu MB allocatable, %dMhz, %uMCU",
- device_id + 1,
- device_name,
- (unsigned int) (device_maxmem_alloc / 1024 / 1024),
- (unsigned int) (device_global_mem / 1024 / 1024),
- (unsigned int) (device_maxclock_frequency),
- (unsigned int) device_processors);
- }
- else
- {
- log_info ("Device #%u: %s, skipped",
- device_id + 1,
- device_name);
+ if (device_param->skipped == 0)
+ {
+ log_info ("Device #%u: %s, %lu/%lu MB allocatable, %dMhz, %uMCU",
+ device_id + 1,
+ device_name,
+ (unsigned int) (device_maxmem_alloc / 1024 / 1024),
+ (unsigned int) (device_global_mem / 1024 / 1024),
+ (unsigned int) (device_maxclock_frequency),
+ (unsigned int) device_processors);
+ }
+ else
+ {
+ log_info ("Device #%u: %s, skipped",
+ device_id + 1,
+ device_name);
+ }
}
}
if ((benchmark == 1 || quiet == 0) && (algorithm_pos == 0))
{
- log_info ("");
+ if (status_automat == 0)
+ {
+ log_info ("");
+ }
}
/**
{
log_info ("Watchdog: Temperature retain trigger set to %uc", gpu_temp_retain);
}
+
+ if (data.quiet == 0) log_info ("");
#endif
}
- if (data.quiet == 0) log_info ("");
-
/**
* HM devices: copy
*/
if (benchmark == 1) log_info ("Hashmode: %d", data.hash_mode);
#endif
+ if (data.quiet == 0) log_info_nn ("Initializing device kernels and memory...");
+
uint kernel_power_all = 0;
for (uint device_id = 0; device_id < data.devices_cnt; device_id++)
device_param->size_root_css = size_root_css;
device_param->size_markov_css = size_markov_css;
- size_t size_results = kernel_threads * sizeof (uint);
+ size_t size_results = sizeof (uint);
device_param->size_results = size_results;
}
}
- if (quiet == 0) log_info ("");
-
for (uint tmto = tmto_start; tmto < tmto_stop; tmto++)
{
// TODO: in theory the following calculation needs to be done per salt, not global
return -1;
}
- if (quiet == 0) log_info ("");
if (quiet == 0) log_info ("SCRYPT tmto optimizer value set to: %u, mem: %u\n", data.salts_buf[0].scrypt_tmto, size_scryptV);
}
* some algorithms have a maximum kernel-loops count
*/
- if (attack_exec == ATTACK_EXEC_OUTSIDE_KERNEL)
+ if (device_param->kernel_loops_min < device_param->kernel_loops_max)
{
- if (data.salts_buf[0].salt_iter < device_param->kernel_loops_max)
+ u32 innerloop_cnt = 0;
+
+ if (data.attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
+ {
+ if (data.attack_kern == ATTACK_KERN_STRAIGHT) innerloop_cnt = data.kernel_rules_cnt;
+ else if (data.attack_kern == ATTACK_KERN_COMBI) innerloop_cnt = data.combs_cnt;
+ else if (data.attack_kern == ATTACK_KERN_BF) innerloop_cnt = data.bfs_cnt;
+ }
+ else
{
- device_param->kernel_loops_max = data.salts_buf[0].salt_iter;
+ innerloop_cnt = data.salts_buf[0].salt_iter;
+ }
+
+ if ((innerloop_cnt >= device_param->kernel_loops_min) &&
+ (innerloop_cnt <= device_param->kernel_loops_max))
+ {
+ device_param->kernel_loops_max = innerloop_cnt;
}
}
case 13000: size_tmps = kernel_power_max * sizeof (pbkdf2_sha256_tmp_t); break;
case 13200: size_tmps = kernel_power_max * sizeof (axcrypt_tmp_t); break;
case 13400: size_tmps = kernel_power_max * sizeof (keepass_tmp_t); break;
+ case 13600: size_tmps = kernel_power_max * sizeof (pbkdf2_sha1_tmp_t); break;
+ case 13711: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13712: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13713: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13721: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break;
+ case 13722: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break;
+ case 13723: size_tmps = kernel_power_max * sizeof (tc64_tmp_t); break;
+ case 13731: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13732: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13733: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13741: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13742: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13743: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13751: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13752: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13753: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13761: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13762: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
+ case 13763: size_tmps = kernel_power_max * sizeof (tc_tmp_t); break;
};
// size_hooks
snprintf (build_opts, sizeof (build_opts) - 1, "-cl-std=CL1.1 -I\"%s/\" -DVENDOR_ID=%u -DCUDA_ARCH=%d -DVECT_SIZE=%u -DDEVICE_TYPE=%u -DKERN_TYPE=%u -D_unroll", shared_dir, device_param->vendor_id, (device_param->sm_major * 100) + device_param->sm_minor, device_param->vector_width, (u32) device_param->device_type, kern_type);
+ if (device_param->vendor_id == VENDOR_ID_INTEL_SDK)
+ {
+ // we do vectorizing much better than the auto-vectorizer
+
+ char build_opts_new[1024] = { 0 };
+
+ snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -cl-opt-disable", build_opts);
+
+ strncpy (build_opts, build_opts_new, sizeof (build_opts) - 1);
+ }
+
+ #ifdef DEBUG
+ log_info ("Device #%u: build_opts '%s'\n", device_id + 1, build_opts);
+ #endif
+
/**
* main kernel
*/
if (rc != 0)
{
device_param->skipped = true;
+
log_info ("Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file);
+
continue;
}
}
else
{
- if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size);
+ #ifdef DEBUG
+ log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size);
+ #endif
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
}
else
{
- if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, source_file, sst.st_size);
+ #ifdef DEBUG
+ log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, source_file, sst.st_size);
+ #endif
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
if (cached == 0)
{
if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, cached_file);
+ if (quiet == 0) log_info ("");
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
if (rc != 0)
{
device_param->skipped = true;
+
log_info ("Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file);
+
continue;
}
}
else
{
- if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size);
+ #ifdef DEBUG
+ log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size);
+ #endif
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
if (cached == 0)
{
if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, cached_file);
+ if (quiet == 0) log_info ("");
load_kernel (source_file, 1, kernel_lengths, kernel_sources);
if (rc != 0)
{
device_param->skipped = true;
+
log_info ("Device #%u: Kernel %s build failure. Proceed without this device.", device_id + 1, source_file);
+
continue;
}
}
else
{
+ #ifdef DEBUG
if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, cached_file, cst.st_size);
+ #endif
load_kernel (cached_file, 1, kernel_lengths, kernel_sources);
* main host data
*/
- uint *result = (uint *) mymalloc (size_results);
-
- device_param->result = result;
-
pw_t *pws_buf = (pw_t *) mymalloc (size_pws);
device_param->pws_buf = pws_buf;
* kernel name
*/
+ size_t kernel_wgs_tmp;
+
char kernel_name[64] = { 0 };
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_tm", kern_type);
device_param->kernel_tm = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
+
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_tm, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
}
}
}
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook12", kern_type);
device_param->kernel12 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
+
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel12, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
}
if (opts_type & OPTS_TYPE_HOOK23)
snprintf (kernel_name, sizeof (kernel_name) - 1, "m%05d_hook23", kern_type);
device_param->kernel23 = hc_clCreateKernel (data.ocl, device_param->program, kernel_name);
+
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel23, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
}
}
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel1, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel2, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel3, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+
for (uint i = 0; i <= 20; i++)
{
hc_clSetKernelArg (data.ocl, device_param->kernel1, i, sizeof (cl_mem), device_param->kernel_params[i]);
device_param->kernel_mp_l = hc_clCreateKernel (data.ocl, device_param->program_mp, "l_markov");
device_param->kernel_mp_r = hc_clCreateKernel (data.ocl, device_param->program_mp, "r_markov");
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp_l, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp_r, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+
if (opts_type & OPTS_TYPE_PT_BITSLICE)
{
hc_clSetKernelArg (data.ocl, device_param->kernel_tm, 0, sizeof (cl_mem), device_param->kernel_params_tm[0]);
else if (attack_mode == ATTACK_MODE_HYBRID1)
{
device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov");
+
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
}
else if (attack_mode == ATTACK_MODE_HYBRID2)
{
device_param->kernel_mp = hc_clCreateKernel (data.ocl, device_param->program_mp, "C_markov");
+
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_mp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
else
{
device_param->kernel_amp = hc_clCreateKernel (data.ocl, device_param->program_amp, "amp");
+
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_amp, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
}
if (attack_exec == ATTACK_EXEC_INSIDE_KERNEL)
}
}
+ // maybe this has been updated by clGetKernelWorkGroupInfo()
+ // value can only be decreased, so we don't need to reallocate buffers
+
+ device_param->kernel_threads = kernel_threads;
+
/**
* Store initial fanspeed if gpu_temp_retain is enabled
*/
data.kernel_power_all = kernel_power_all;
- if (data.quiet == 0) log_info ("");
+ if (data.quiet == 0) log_info_nn ("");
/**
* In benchmark-mode, inform user which algorithm is checked
if (benchmark == 1)
{
- quiet = 0;
+ if (status_automat == 0)
+ {
+ quiet = 0;
- data.quiet = quiet;
+ data.quiet = quiet;
- char *hash_type = strhashtype (data.hash_mode); // not a bug
+ char *hash_type = strhashtype (data.hash_mode); // not a bug
- log_info ("Hashtype: %s", hash_type);
- log_info ("");
+ log_info ("Hashtype: %s", hash_type);
+ log_info ("");
+ }
}
/**
{
weak_hash_check (device_param, salt_pos);
}
- }
- // Display hack, guarantee that there is at least one \r before real start
+ // Display hack, guarantee that there is at least one \r before real start
- if (data.quiet == 0) log_info_nn ("");
+ //if (data.quiet == 0) log_info ("");
+ }
/**
* status and monitor threads
if (data.outfile_check_directory != NULL)
{
if ((hash_mode != 5200) &&
- !((hash_mode >= 6200) && (hash_mode <= 6299)) &&
+ !((hash_mode >= 6200) && (hash_mode <= 6299)) &&
+ !((hash_mode >= 13700) && (hash_mode <= 13799)) &&
(hash_mode != 9000))
{
hc_thread_create (ni_threads[ni_threads_cnt], thread_outfile_remove, NULL);
{
if (quiet == 0)
{
- log_info ("");
log_info ("ATTENTION!");
log_info (" The wordlist or mask you are using is too small.");
log_info (" Therefore, hashcat is unable to utilize the full parallelization power of your device(s).");
{
status_benchmark ();
- log_info ("");
+ if (status_automat == 0)
+ {
+ log_info ("");
+ }
}
else
{
if (device_param->skipped) continue;
- local_free (device_param->result);
-
local_free (device_param->combs_buf);
local_free (device_param->hooks_buf);
projects / hashcat.git / blobdiff