const uint VERSION_BIN = 300;
const uint RESTORE_MIN = 300;
-double TARGET_MS_PROFILE[4] = { 8, 16, 96, 512 };
+double TARGET_MS_PROFILE[4] = { 2, 12, 96, 480 };
#define INCR_RULES 10000
#define INCR_SALTS 100000
#define RESTORE_DISABLE 0
#define STATUS 0
#define STATUS_TIMER 10
-#define STATUS_AUTOMAT 0
+#define MACHINE_READABLE 0
#define LOOPBACK 0
#define WEAK_HASH_THRESHOLD 100
#define SHOW 0
#define BITMAP_MAX 24
#define GPU_TEMP_DISABLE 0
#define GPU_TEMP_ABORT 90
-#define GPU_TEMP_RETAIN 80
+#define GPU_TEMP_RETAIN 70
#define WORKLOAD_PROFILE 2
#define KERNEL_ACCEL 0
#define KERNEL_LOOPS 0
"",
"Usage: %s [options]... hash|hashfile|hccapfile [dictionary|mask|directory]...",
"",
- "###########",
- "# Options #",
- "###########",
+ "- [ Options ] -",
"",
- " Short- | Long- | Parameter- | Description | Example",
- " Option | Option | Type | |",
- "--------+---------------------------+------------+------------------------------------------------------+---------------------",
- " -m | --hash-type | Number | Hash-type, see references below | -m 1000",
- " -a | --attack-mode | Number | Attack-mode, see references below | -a 3",
- " -V | --version | | Print version | --version",
- " -h | --help | | Print help | --help",
- " | --quiet | | Suppress output | --quiet",
- " | --hex-charset | | Assume charset is given in hex | --hex-charset",
- " | --hex-salt | | Assume salt is given in hex | --hex-salt",
- " | --hex-wordlist | | Assume words in wordlist is given in hex | --hex-wordlist",
- " | --force | | Ignore warnings | --force",
- " | --status | | Enable automatic update of the status-screen | --status",
- " | --status-timer | Number | Sets seconds between status-screen update to X | --status-timer=1",
- " | --status-automat | | Display the status view in a machine readable format | --status-automat",
- " | --loopback | | Add new plains to induct directory | --loopback",
- " | --weak-hash-threshold | Number | Threshold X when to stop checking for weak hashes | --weak-hash-threshold=0",
- " | --markov-hcstat | File | Specify hcstat file to use, default is hashcat.hcstat| --markiv-hcstat=my.hcstat",
- " | --markov-disable | | Disables markov-chains, emulates classic brute-force | --markov-disable",
- " | --markov-classic | | Enables classic markov-chains, no per-position | --markov-classic",
- " -t | --markov-threshold | Number | Threshold X when to stop accepting new markov-chains | -t 50",
- " | --runtime | Number | Abort session after X seconds of runtime | --runtime=10",
- " | --session | String | Define specific session name | --session=mysession",
- " | --restore | | Restore session from --session | --restore",
- " | --restore-disable | | Do not write restore file | --restore-disable",
- " -o | --outfile | File | Define outfile for recovered hash | -o outfile.txt",
- " | --outfile-format | Number | Define outfile-format X for recovered hash | --outfile-format=7",
- " | --outfile-autohex-disable | | Disable the use of $HEX[] in output plains | --outfile-autohex-disable",
- " | --outfile-check-timer | Number | Sets seconds between outfile checks to X | --outfile-check-timer=30",
- " -p | --separator | Char | Separator char for hashlists and outfile | -p :",
- " | --show | | Show cracked passwords only | --show",
- " | --left | | Show un-cracked passwords only | --left",
- " | --username | | Enable ignoring of usernames in hashfile | --username",
- " | --remove | | Enable remove of hash once it is cracked | --remove",
- " | --remove-timer | Number | Update input hash file each X seconds | --remove-timer=30",
- " | --potfile-disable | | Do not write potfile | --potfile-disable",
- " | --potfile-path | | Specific path to potfile | --potfile-path=hashcat.pot",
- " | --debug-mode | Number | Defines the debug mode (hybrid only by using rules) | --debug-mode=4",
- " | --debug-file | File | Output file for debugging rules | --debug-file=goodrules.log",
- " | --induction-dir | Folder | Specify the induction directory to use for loopback | --induction-dir=inducts",
- " | --outfile-check-dir | Folder | Specify the outfile directory to monitor for plains | --outfile-check-dir=ext",
- " | --logfile-disable | | Disable the logfile | --logfile-disable",
- " | --truecrypt-keyfiles | File | Keyfiles used, separate with comma | --truecrypt-keyfiles=xx.png",
- " | --veracrypt-keyfiles | File | Keyfiles used, separate with comma | --veracrypt-keyfiles=xx.png",
- " | --veracrypt-pim | Number | VeraCrypt personal iterations multiplier | --veracrypt-pim=1000",
- " -b | --benchmark | | Run benchmark | -b",
- " -c | --segment-size | Number | Sets size in MB to cache from the wordfile to X | -c 32",
- " | --bitmap-min | Number | Sets minimum bits allowed for bitmaps to X | --bitmap-min=24",
- " | --bitmap-max | Number | Sets maximum bits allowed for bitmaps to X | --bitmap-min=24",
- " | --cpu-affinity | String | Locks to CPU devices, separate with comma | --cpu-affinity=1,2,3,4",
- " | --opencl-platforms | String | OpenCL platforms to use, separate with comma | --opencl-platforms=2",
- " -d | --opencl-devices | String | OpenCL devices to use, separate with comma | -d 1",
- " | --opencl-device-types | String | OpenCL device-types to use, separate with comma | --opencl-device-types=1,2",
- " | --opencl-vector-width | Number | Manual override OpenCL vector-width to X | --opencl-vector-width=4",
- " -w | --workload-profile | Number | Enable a specific workload profile, see pool below | -w 3",
- " -n | --kernel-accel | Number | Manual workload tuning, set outerloop step size to X | -n 64",
- " -u | --kernel-loops | Number | Manual workload tuning, set innerloop step size to X | -u 256",
- " | --gpu-temp-disable | | Disable temperature and fanspeed reads and triggers | --gpu-temp-disable",
+ " Options Short / Long | Type | Description | Example",
+ "===============================|======|======================================================|=======================",
+ " -m, --hash-type | Num | Hash-type, see references below | -m 1000",
+ " -a, --attack-mode | Num | Attack-mode, see references below | -a 3",
+ " -V, --version | | Print version |",
+ " -h, --help | | Print help |",
+ " --quiet | | Suppress output |",
+ " --hex-charset | | Assume charset is given in hex |",
+ " --hex-salt | | Assume salt is given in hex |",
+ " --hex-wordlist | | Assume words in wordlist is given in hex |",
+ " --force | | Ignore warnings |",
+ " --status | | Enable automatic update of the status-screen |",
+ " --status-timer | Num | Sets seconds between status-screen update to X | --status-timer=1",
+ " --machine-readable | | Display the status view in a machine readable format |",
+ " --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",
+ " --markov-disable | | Disables markov-chains, emulates classic brute-force |",
+ " --markov-classic | | Enables classic markov-chains, no per-position |",
+ " -t, --markov-threshold | Num | Threshold X when to stop accepting new markov-chains | -t 50",
+ " --runtime | Num | Abort session after X seconds of runtime | --runtime=10",
+ " --session | Str | Define specific session name | --session=mysession",
+ " --restore | | Restore session from --session |",
+ " --restore-disable | | Do not write restore file |",
+ " -o, --outfile | File | Define outfile for recovered hash | -o outfile.txt",
+ " --outfile-format | Num | Define outfile-format X for recovered hash | --outfile-format=7",
+ " --outfile-autohex-disable | | Disable the use of $HEX[] in output plains |",
+ " --outfile-check-timer | Num | Sets seconds between outfile checks to X | --outfile-check=30",
+ " -p, --separator | Char | Separator char for hashlists and outfile | -p :",
+ " --show | | Show cracked passwords only |",
+ " --left | | Show un-cracked passwords only |",
+ " --username | | Enable ignoring of usernames in hashfile |",
+ " --remove | | Enable remove of hash once it is cracked |",
+ " --remove-timer | Num | Update input hash file each X seconds | --remove-timer=30",
+ " --potfile-disable | | Do not write potfile |",
+ " --potfile-path | Dir | Specific path to potfile | --potfile-path=my.pot",
+ " --debug-mode | Num | Defines the debug mode (hybrid only by using rules) | --debug-mode=4",
+ " --debug-file | File | Output file for debugging rules | --debug-file=good.log",
+ " --induction-dir | Dir | Specify the induction directory to use for loopback | --induction=inducts",
+ " --outfile-check-dir | Dir | Specify the outfile directory to monitor for plains | --outfile-check-dir=x",
+ " --logfile-disable | | Disable the logfile |",
+ " --truecrypt-keyfiles | File | Keyfiles used, separate with comma | --truecrypt-key=x.png",
+ " --veracrypt-keyfiles | File | Keyfiles used, separate with comma | --veracrypt-key=x.txt",
+ " --veracrypt-pim | Num | VeraCrypt personal iterations multiplier | --veracrypt-pim=1000",
+ " -b, --benchmark | | Run benchmark |",
+ " -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-min=24",
+ " --cpu-affinity | Str | Locks to CPU devices, separate with comma | --cpu-affinity=1,2,3",
+ " --opencl-platforms | Str | OpenCL platforms to use, separate with comma | --opencl-platforms=2",
+ " -d, --opencl-devices | Str | OpenCL devices to use, separate with comma | -d 1",
+ " --opencl-device-types | Str | OpenCL device-types to use, separate with comma | --opencl-device-type=1",
+ " --opencl-vector-width | Num | Manual override OpenCL vector-width to X | --opencl-vector=4",
+ " -w, --workload-profile | Num | Enable a specific workload profile, see pool below | -w 3",
+ " -n, --kernel-accel | Num | Manual workload tuning, set outerloop step size to X | -n 64",
+ " -u, --kernel-loops | Num | Manual workload tuning, set innerloop step size to X | -u 256",
+ " --gpu-temp-disable | | Disable temperature and fanspeed reads and triggers |",
#ifdef HAVE_HWMON
- " | --gpu-temp-abort | Number | Abort if GPU temperature reaches X degrees celsius | --gpu-temp-abort=100",
- " | --gpu-temp-retain | Number | Try to retain GPU temperature at X degrees celsius | --gpu-temp-retain=95",
+ " --gpu-temp-abort | Num | Abort if GPU temperature reaches X degrees celsius | --gpu-temp-abort=100",
+ " --gpu-temp-retain | Num | Try to retain GPU temperature at X degrees celsius | --gpu-temp-retain=95",
#ifdef HAVE_ADL
- " | --powertune-enable | | Enable automatic power tuning (AMD OverDrive 6 only) | --powertune-enable",
+ " --powertune-enable | | Enable automatic power tuning (AMD OverDrive 6 only) |",
#endif
#endif
- " | --scrypt-tmto | Number | Manually override TMTO value for scrypt to X | --scrypt-tmto=3",
- " -s | --skip | Number | Skip X words from the start | -s 1000000",
- " -l | --limit | Number | Limit X words from the start + skipped words | -l 1000000",
- " | --keyspace | | Show keyspace base:mod values and quit | --keyspace",
- " -j | --rule-left | Rule | Single Rule applied to each word from left wordlist | -j 'c'",
- " -k | --rule-right | Rule | Single Rule applied to each word from right wordlist | -k '^-'",
- " -r | --rules-file | File | Multiple Rules applied to each word from wordlists | -r rules/best64.rule",
- " -g | --generate-rules | | Enable generate random rules | -g 10000",
- " | --generate-rules-func-min | Number | Force min X funcs per rule |",
- " | --generate-rules-func-max | Number | Force max X funcs per rule |",
- " | --generate-rules-seed | Number | Force RNG seed set to X |",
- " -1 | --custom-charset1 | Charset | User-defined charset ?1 | -1 ?l?d?u",
- " -2 | --custom-charset2 | Charset | User-defined charset ?2 | -2 ?l?d?s",
- " -3 | --custom-charset3 | Charset | User-defined charset ?3 |",
- " -4 | --custom-charset4 | Charset | User-defined charset ?4 |",
- " -i | --increment | | Enable mask increment mode | --increment",
- " | --increment-min | Number | Start mask incrementing at X | --increment-min=4",
- " | --increment-max | Number | Stop mask incrementing at X | --increment-max=8",
+ " --scrypt-tmto | Num | Manually override TMTO value for scrypt to X | --scrypt-tmto=3",
+ " -s, --skip | Num | Skip X words from the start | -s 1000000",
+ " -l, --limit | Num | Limit X words from the start + skipped words | -l 1000000",
+ " --keyspace | | Show keyspace base:mod values and quit |",
+ " -j, --rule-left | Rule | Single Rule applied to each word from left wordlist | -j 'c'",
+ " -k, --rule-right | Rule | Single Rule applied to each word from right wordlist | -k '^-'",
+ " -r, --rules-file | File | Multiple Rules applied to each word from wordlists | -r rules/best64.rule",
+ " -g, --generate-rules | Num | Generate X random rules | -g 10000",
+ " --generate-rules-func-min | Num | Force min X funcs per rule |",
+ " --generate-rules-func-max | Num | Force max X funcs per rule |",
+ " --generate-rules-seed | Num | Force RNG seed set to X |",
+ " -1, --custom-charset1 | CS | User-defined charset ?1 | -1 ?l?d?u",
+ " -2, --custom-charset2 | CS | User-defined charset ?2 | -2 ?l?d?s",
+ " -3, --custom-charset3 | CS | User-defined charset ?3 |",
+ " -4, --custom-charset4 | CS | User-defined charset ?4 |",
+ " -i, --increment | | Enable mask increment mode |",
+ " --increment-min | Num | Start mask incrementing at X | --increment-min=4",
+ " --increment-max | Num | Stop mask incrementing at X | --increment-max=8",
"",
- "################",
- "## Hash modes ##",
- "################",
+ "- [ Hash modes ] -",
"",
" # | Name | Category",
- " ------+--------------------------------------------------+--------------------------------------",
+ " ======+==================================================+======================================",
" 900 | MD4 | Raw Hash",
" 0 | MD5 | Raw Hash",
" 5100 | Half MD5 | Raw Hash",
" 12700 | Blockchain, My Wallet | Password Managers",
" 13400 | Keepass 1 (AES/Twofish) and Keepass 2 (AES) | Password Managers",
"",
- "#####################",
- "## Outfile Formats ##",
- "#####################",
+ "- [ Outfile Formats ] -",
"",
" # | Format",
- " ---+--------",
+ " ===+========",
" 1 | hash[:salt]",
" 2 | plain",
" 3 | hash[:salt]:plain",
" 14 | plain:hex_plain:crack_pos",
" 15 | hash[:salt]:plain:hex_plain:crack_pos",
"",
- "##########################",
- "## Rule Debugging Modes ##",
- "##########################",
+ "- [ Rule Debugging Modes ] -",
"",
" # | Format",
- " ---+--------",
+ " ===+========",
" 1 | Finding-Rule",
" 2 | Original-Word",
" 3 | Original-Word:Finding-Rule",
" 4 | Original-Word:Finding-Rule:Processed-Word",
"",
- "##################",
- "## Attack Modes ##",
- "##################",
+ "- [ Attack Modes ] -",
"",
" # | Mode",
- " ---+------",
+ " ===+======",
" 0 | Straight",
" 1 | Combination",
" 3 | Brute-force",
" 6 | Hybrid Wordlist + Mask",
" 7 | Hybrid Mask + Wordlist",
"",
- "#######################",
- "## Built-in Charsets ##",
- "#######################",
+ "- [ Built-in Charsets ] -",
"",
" ? | Charset",
- " ---+---------",
+ " ===+=========",
" l | abcdefghijklmnopqrstuvwxyz",
" u | ABCDEFGHIJKLMNOPQRSTUVWXYZ",
" d | 0123456789",
" a | ?l?u?d?s",
" b | 0x00 - 0xff",
"",
- "#########################",
- "## OpenCL Device Types ##",
- "#########################",
+ "- [ OpenCL Device Types ] -",
"",
" # | Device Type",
- " ---+-------------",
+ " ===+=============",
" 1 | CPU",
" 2 | GPU",
" 3 | FPGA, DSP, Co-Processor",
"",
- "#######################",
- "## Workload Profiles ##",
- "#######################",
+ "- [ Workload Profiles ] -",
"",
" # | Performance | Runtime | Power Consumption | Desktop Impact",
- " ---+-------------+---------+-------------------+----------------",
- " 1 | Low | 8 ms | Low | Minimal",
- " 2 | Default | 16 ms | Economic | Noticeable",
+ " ===+=============+=========+===================+=================",
+ " 1 | Low | 2 ms | Low | Minimal",
+ " 2 | Default | 12 ms | Economic | Noticeable",
" 3 | High | 96 ms | High | Unresponsive",
- " 4 | Nightmare | 512 ms | Insane | Headless",
+ " 4 | Nightmare | 480 ms | Insane | Headless",
"",
- "If you have no idea what just happend, please visit the following pages:",
+ "If you have no idea what just happened then visit the following pages:",
"",
"* https://hashcat.net/wiki/#howtos_videos_papers_articles_etc_in_the_wild",
"* https://hashcat.net/wiki/#frequently_asked_questions",
return exec_ms_sum / exec_ms_cnt;
}
-void status_display_automat ()
+void status_display_machine_readable ()
{
FILE *out = stdout;
if (data.devices_status == STATUS_STARTING) return;
if (data.devices_status == STATUS_BYPASS) return;
- if (data.status_automat == 1)
+ if (data.machine_readable == 1)
{
- status_display_automat ();
+ status_display_machine_readable ();
return;
}
if (device_param->skipped) continue;
- #define HM_STR_BUF_SIZE 255
+ const int num_temperature = hm_get_temperature_with_device_id (device_id);
+ const int num_fanspeed = hm_get_fanspeed_with_device_id (device_id);
+ const int num_utilization = hm_get_utilization_with_device_id (device_id);
+ const int num_corespeed = hm_get_corespeed_with_device_id (device_id);
+ const int num_memoryspeed = hm_get_memoryspeed_with_device_id (device_id);
+ const int num_buslanes = hm_get_buslanes_with_device_id (device_id);
+ // not working
+ //const int num_throttle = hm_get_throttle_with_device_id (device_id);
- if (data.hm_device[device_id].fan_supported == 1)
+ char output_buf[256] = { 0 };
+
+ int output_len = 0;
+
+ if (num_temperature >= 0)
{
- char utilization[HM_STR_BUF_SIZE] = { 0 };
- char temperature[HM_STR_BUF_SIZE] = { 0 };
- char fanspeed[HM_STR_BUF_SIZE] = { 0 };
+ snprintf (output_buf + output_len, sizeof (output_buf) - output_len, " Temp:%3uc", num_temperature);
- hm_device_val_to_str ((char *) utilization, HM_STR_BUF_SIZE, "%", hm_get_utilization_with_device_id (device_id));
- hm_device_val_to_str ((char *) temperature, HM_STR_BUF_SIZE, "c", hm_get_temperature_with_device_id (device_id));
+ output_len = strlen (output_buf);
+ }
- if (device_param->vendor_id == VENDOR_ID_AMD)
- {
- hm_device_val_to_str ((char *) fanspeed, HM_STR_BUF_SIZE, "%", hm_get_fanspeed_with_device_id (device_id));
- }
- else if (device_param->vendor_id == VENDOR_ID_NV)
- {
- hm_device_val_to_str ((char *) fanspeed, HM_STR_BUF_SIZE, "%", hm_get_fanspeed_with_device_id (device_id));
- }
+ if (num_fanspeed >= 0)
+ {
+ snprintf (output_buf + output_len, sizeof (output_buf) - output_len, " Fan:%3u%%", num_fanspeed);
- log_info ("HWMon.GPU.#%d...: %s Util, %s Temp, %s Fan", device_id + 1, utilization, temperature, fanspeed);
+ output_len = strlen (output_buf);
}
- else
+
+ if (num_utilization >= 0)
+ {
+ snprintf (output_buf + output_len, sizeof (output_buf) - output_len, " Util:%3u%%", num_utilization);
+
+ output_len = strlen (output_buf);
+ }
+
+ if (num_corespeed >= 0)
+ {
+ snprintf (output_buf + output_len, sizeof (output_buf) - output_len, " Core:%4uMhz", num_corespeed);
+
+ output_len = strlen (output_buf);
+ }
+
+ if (num_memoryspeed >= 0)
{
- char utilization[HM_STR_BUF_SIZE] = { 0 };
- char temperature[HM_STR_BUF_SIZE] = { 0 };
+ snprintf (output_buf + output_len, sizeof (output_buf) - output_len, " Mem:%4uMhz", num_memoryspeed);
- hm_device_val_to_str ((char *) utilization, HM_STR_BUF_SIZE, "%", hm_get_utilization_with_device_id (device_id));
- hm_device_val_to_str ((char *) temperature, HM_STR_BUF_SIZE, "c", hm_get_temperature_with_device_id (device_id));
+ output_len = strlen (output_buf);
+ }
+
+ if (num_buslanes >= 0)
+ {
+ snprintf (output_buf + output_len, sizeof (output_buf) - output_len, " Lanes:%u", num_buslanes);
+
+ output_len = strlen (output_buf);
+ }
+
+ /*
+ if (num_throttle >= 0)
+ {
+ snprintf (output_buf + output_len, sizeof (output_buf) - output_len, " Throttle:%u", num_throttle);
+
+ output_len = strlen (output_buf);
+ }
+ */
- log_info ("HWMon.GPU.#%d...: %s Util, %s Temp, N/A Fan", device_id + 1, utilization, temperature);
+ if (output_len == 0)
+ {
+ snprintf (output_buf + output_len, sizeof (output_buf) - output_len, " N/A");
+
+ output_len = strlen (output_buf);
}
+
+ log_info ("HWMon.Dev.#%d...:%s", device_id + 1, output_buf);
}
hc_thread_mutex_unlock (mux_adl);
#endif // HAVE_HWMON
}
-static void status_benchmark_automat ()
+static void status_benchmark_automate ()
{
u64 speed_cnt[DEVICES_MAX] = { 0 };
double speed_ms[DEVICES_MAX] = { 0 };
if (data.devices_status == STATUS_STARTING) return;
if (data.devices_status == STATUS_BYPASS) return;
- if (data.status_automat == 1)
+ if (data.machine_readable == 1)
{
- status_benchmark_automat ();
+ status_benchmark_automate ();
return;
}
out_fp = stdout;
}
+
lock_file (out_fp);
}
else
if (data.digests_shown[hash_pos] == 1) continue;
+ hc_thread_mutex_lock (mux_display);
+
if ((data.opts_type & OPTS_TYPE_PT_NEVERCRACK) == 0)
{
data.digests_shown[hash_pos] = 1;
if (data.salts_done == data.salts_cnt) data.devices_status = STATUS_CRACKED;
+ hc_thread_mutex_unlock (mux_display);
+
check_hash (device_param, &cracked[i]);
}
if (cpt_cracked > 0)
{
+ hc_thread_mutex_lock (mux_display);
+
data.cpt_buf[data.cpt_pos].timestamp = time (NULL);
data.cpt_buf[data.cpt_pos].cracked = cpt_cracked;
data.cpt_total += cpt_cracked;
if (data.cpt_pos == CPT_BUF) data.cpt_pos = 0;
+
+ hc_thread_mutex_unlock (mux_display);
}
if (data.opts_type & OPTS_TYPE_PT_NEVERCRACK)
hc_clFinish (data.ocl, device_param->command_queue);
}
+static void run_kernel_memset (hc_device_param_t *device_param, cl_mem buf, const uint value, const uint num)
+{
+ const u32 num16d = num / 16;
+ const u32 num16m = num % 16;
+
+ if (num16d)
+ {
+ device_param->kernel_params_memset_buf32[1] = value;
+ device_param->kernel_params_memset_buf32[2] = num16d;
+
+ uint kernel_threads = device_param->kernel_threads;
+
+ uint num_elements = num16d;
+
+ while (num_elements % kernel_threads) num_elements++;
+
+ cl_kernel kernel = device_param->kernel_memset;
+
+ hc_clSetKernelArg (data.ocl, kernel, 0, sizeof (cl_mem), (void *) &buf);
+ hc_clSetKernelArg (data.ocl, kernel, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]);
+ hc_clSetKernelArg (data.ocl, kernel, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]);
+
+ 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_clFlush (data.ocl, device_param->command_queue);
+
+ hc_clFinish (data.ocl, device_param->command_queue);
+ }
+
+ if (num16m)
+ {
+ u32 tmp[4];
+
+ tmp[0] = value;
+ tmp[1] = value;
+ tmp[2] = value;
+ tmp[3] = value;
+
+ hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, buf, CL_TRUE, num16d * 16, num16m, tmp, 0, NULL, NULL);
+ }
+}
+
static void run_kernel_bzero (hc_device_param_t *device_param, cl_mem buf, const size_t size)
{
+ run_kernel_memset (device_param, buf, 0, size);
+
+ /*
int rc = -1;
- if (device_param->opencl_v12 && device_param->vendor_id == VENDOR_ID_AMD)
+ if (device_param->opencl_v12 && device_param->platform_vendor_id == VENDOR_ID_AMD)
{
// So far tested, amd is the only supporting this OpenCL 1.2 function without segfaulting
myfree (tmp);
}
+ */
}
static void choose_kernel (hc_device_param_t *device_param, const uint attack_exec, const uint attack_mode, const uint opts_type, const salt_t *salt_buf, const uint highest_pw_len, const uint pws_cnt)
}
else if (data.attack_kern == ATTACK_KERN_COMBI)
{
- if (data.attack_mode == ATTACK_MODE_HYBRID2)
+ if (data.attack_mode == ATTACK_MODE_COMBI)
+ {
+ if (data.combs_mode == COMBINATOR_MODE_BASE_RIGHT)
+ {
+ if (data.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 (data.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 (data.attack_mode == ATTACK_MODE_HYBRID2)
{
if (data.opts_type & OPTS_TYPE_PT_ADD01)
{
try_run (device_param, kernel_accel, kernel_loops);
try_run (device_param, kernel_accel, kernel_loops);
+ device_param->kernel_accel = kernel_accel;
+ device_param->kernel_loops = kernel_loops;
+
+ const u32 kernel_power = device_param->device_processors * device_param->kernel_threads * device_param->kernel_accel;
+
+ device_param->kernel_power = kernel_power;
+
return;
}
const u32 kernel_power_max = device_param->device_processors * device_param->kernel_threads * kernel_accel_max;
+ run_kernel_memset (device_param, device_param->d_pws_buf, 7, kernel_power_max * sizeof (pw_t));
+
+ if (data.attack_exec == ATTACK_EXEC_OUTSIDE_KERNEL)
+ {
+ run_kernel_memset (device_param, device_param->d_pws_amp_buf, 7, kernel_power_max * sizeof (pw_t));
+ }
+
+ /*
for (u32 i = 0; i < kernel_power_max; i++)
{
device_param->pws_buf[i].i[0] = i;
{
run_kernel_amp (device_param, kernel_power_max);
}
+ */
#define VERIFIER_CNT 1
exec_ms_pre_final = MIN (exec_ms_pre_final, exec_ms_pre_final_v);
}
+ u32 diff = kernel_loops - kernel_accel;
+
if ((kernel_loops_min < kernel_loops_max) && (kernel_accel_min < kernel_accel_max))
{
- for (u32 f = 2; f < 1024; f++)
+ u32 kernel_accel_orig = kernel_accel;
+ u32 kernel_loops_orig = kernel_loops;
+
+ for (u32 f = 1; f < 1024; f++)
{
- const u32 kernel_accel_try = kernel_accel * f;
- const u32 kernel_loops_try = kernel_loops / f;
+ const u32 kernel_accel_try = (float) kernel_accel_orig * f;
+ const u32 kernel_loops_try = (float) kernel_loops_orig / f;
if (kernel_accel_try > kernel_accel_max) break;
if (kernel_loops_try < kernel_loops_min) break;
+ u32 diff_new = kernel_loops_try - kernel_accel_try;
+
+ if (diff_new > diff) break;
+
+ diff_new = diff;
+
double exec_ms = try_run (device_param, kernel_accel_try, kernel_loops_try);
for (int i = 0; i < VERIFIER_CNT; i++)
const double accel_left = kernel_accel_max / kernel_accel;
- const int exec_accel_min = MIN (exec_left, accel_left); // we want that to be int
+ const double exec_accel_min = MIN (exec_left, accel_left); // we want that to be int
- if (exec_accel_min >= 2)
+ if (exec_accel_min >= 1.0)
{
// this is safe to not overflow kernel_accel_max because of accel_left
- kernel_accel *= exec_accel_min;
+ kernel_accel = (double) kernel_accel * exec_accel_min;
}
// reset them fake words
+ /*
memset (device_param->pws_buf, 0, kernel_power_max * sizeof (pw_t));
hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_buf, CL_TRUE, 0, kernel_power_max * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_pws_amp_buf, CL_TRUE, 0, kernel_power_max * sizeof (pw_t), device_param->pws_buf, 0, NULL, NULL);
+ */
+
+ run_kernel_memset (device_param, device_param->d_pws_buf, 0, kernel_power_max * sizeof (pw_t));
+
+ if (data.attack_exec == ATTACK_EXEC_OUTSIDE_KERNEL)
+ {
+ run_kernel_memset (device_param, device_param->d_pws_amp_buf, 0, kernel_power_max * sizeof (pw_t));
+ }
// reset timer
* result
*/
- hc_thread_mutex_lock (mux_display);
-
check_cracked (device_param, salt_pos);
- hc_thread_mutex_unlock (mux_display);
-
/**
* progress
*/
hc_timer_set (&device_param->timer_speed);
- hc_thread_mutex_lock (mux_display);
-
// current speed
+ //hc_thread_mutex_lock (mux_display);
+
device_param->speed_cnt[speed_pos] = perf_sum_all;
device_param->speed_ms[speed_pos] = speed_ms;
- hc_thread_mutex_unlock (mux_display);
+ //hc_thread_mutex_unlock (mux_display);
speed_pos++;
#ifdef HAVE_HWMON
uint hwmon_check = 0;
+ int slowdown_warnings = 0;
+
// these variables are mainly used for fan control (AMD only)
int *fan_speed_chgd = (int *) mycalloc (data.devices_cnt, sizeof (int));
if (data.devices_status != STATUS_RUNNING) continue;
#ifdef HAVE_HWMON
+
+ if (1)
+ {
+ hc_thread_mutex_lock (mux_adl);
+
+ 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;
+
+ if ((data.devices_param[device_id].device_type & CL_DEVICE_TYPE_GPU) == 0) continue;
+
+ const int temperature = hm_get_temperature_with_device_id (device_id);
+
+ const int threshold = data.hm_device[device_id].gpu_temp_threshold_slowdown;
+
+ if (temperature >= threshold)
+ {
+ if (slowdown_warnings < 3)
+ {
+ if (data.quiet == 0) clear_prompt ();
+
+ log_info ("WARNING: Drivers temperature threshold (%dc) hit on GPU #%d, expect performance to drop...", threshold, device_id + 1);
+
+ if (slowdown_warnings == 2)
+ {
+ log_info ("");
+ }
+
+ if (data.quiet == 0) fprintf (stdout, "%s", PROMPT);
+ if (data.quiet == 0) fflush (stdout);
+
+ slowdown_warnings++;
+ }
+ }
+ else
+ {
+ slowdown_warnings = 0;
+ }
+ }
+
+ hc_thread_mutex_unlock (mux_adl);
+ }
+
if (hwmon_check == 1)
{
hc_thread_mutex_lock (mux_adl);
if (status_left == 0)
{
- hc_thread_mutex_lock (mux_display);
+ //hc_thread_mutex_lock (mux_display);
if (data.quiet == 0) clear_prompt ();
if (data.quiet == 0) log_info ("");
- hc_thread_mutex_unlock (mux_display);
+ //hc_thread_mutex_unlock (mux_display);
status_left = data.status_timer;
}
* main
*/
+#ifdef _WIN
+void SetConsoleWindowSize (const int x)
+{
+ HANDLE h = GetStdHandle (STD_OUTPUT_HANDLE);
+
+ if (h == INVALID_HANDLE_VALUE) return;
+
+ CONSOLE_SCREEN_BUFFER_INFO bufferInfo;
+
+ if (!GetConsoleScreenBufferInfo (h, &bufferInfo)) return;
+
+ SMALL_RECT *sr = &bufferInfo.srWindow;
+
+ sr->Right = MAX (sr->Right, x - 1);
+
+ COORD co;
+
+ co.X = sr->Right + 1;
+ co.Y = sr->Bottom + 1;
+
+ if (!SetConsoleScreenBufferSize (h, co)) return;
+
+ if (!SetConsoleWindowInfo (h, TRUE, sr)) return;
+}
+#endif
+
int main (int argc, char **argv)
{
+ #ifdef _WIN
+ SetConsoleWindowSize (132);
+ #endif
+
/**
* To help users a bit
*/
* commandline parameters
*/
- uint usage = USAGE;
- uint version = VERSION;
- uint quiet = QUIET;
- uint benchmark = BENCHMARK;
- uint show = SHOW;
- uint left = LEFT;
- uint username = USERNAME;
- uint remove = REMOVE;
- uint remove_timer = REMOVE_TIMER;
- u64 skip = SKIP;
- u64 limit = LIMIT;
- uint keyspace = KEYSPACE;
- uint potfile_disable = POTFILE_DISABLE;
- char *potfile_path = NULL;
- uint debug_mode = DEBUG_MODE;
- char *debug_file = NULL;
- char *induction_dir = NULL;
- char *outfile_check_dir = NULL;
- uint force = FORCE;
- uint runtime = RUNTIME;
- uint hash_mode = HASH_MODE;
- uint attack_mode = ATTACK_MODE;
- uint markov_disable = MARKOV_DISABLE;
- uint markov_classic = MARKOV_CLASSIC;
- uint markov_threshold = MARKOV_THRESHOLD;
- char *markov_hcstat = NULL;
- char *outfile = NULL;
- uint outfile_format = OUTFILE_FORMAT;
- uint outfile_autohex = OUTFILE_AUTOHEX;
- uint outfile_check_timer = OUTFILE_CHECK_TIMER;
- uint restore = RESTORE;
- uint restore_timer = RESTORE_TIMER;
- uint restore_disable = RESTORE_DISABLE;
- uint status = STATUS;
- uint status_timer = STATUS_TIMER;
- uint status_automat = STATUS_AUTOMAT;
- uint loopback = LOOPBACK;
- uint weak_hash_threshold = WEAK_HASH_THRESHOLD;
- char *session = NULL;
- uint hex_charset = HEX_CHARSET;
- uint hex_salt = HEX_SALT;
- uint hex_wordlist = HEX_WORDLIST;
- uint rp_gen = RP_GEN;
- uint rp_gen_func_min = RP_GEN_FUNC_MIN;
- uint rp_gen_func_max = RP_GEN_FUNC_MAX;
- uint rp_gen_seed = RP_GEN_SEED;
- char *rule_buf_l = (char *) RULE_BUF_L;
- char *rule_buf_r = (char *) RULE_BUF_R;
- uint increment = INCREMENT;
- uint increment_min = INCREMENT_MIN;
- uint increment_max = INCREMENT_MAX;
- char *cpu_affinity = NULL;
- OCL_PTR *ocl = NULL;
- char *opencl_devices = NULL;
- char *opencl_platforms = NULL;
- 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;
- uint gpu_temp_disable = GPU_TEMP_DISABLE;
+ uint usage = USAGE;
+ uint version = VERSION;
+ uint quiet = QUIET;
+ uint benchmark = BENCHMARK;
+ uint show = SHOW;
+ uint left = LEFT;
+ uint username = USERNAME;
+ uint remove = REMOVE;
+ uint remove_timer = REMOVE_TIMER;
+ u64 skip = SKIP;
+ u64 limit = LIMIT;
+ uint keyspace = KEYSPACE;
+ uint potfile_disable = POTFILE_DISABLE;
+ char *potfile_path = NULL;
+ uint debug_mode = DEBUG_MODE;
+ char *debug_file = NULL;
+ char *induction_dir = NULL;
+ char *outfile_check_dir = NULL;
+ uint force = FORCE;
+ uint runtime = RUNTIME;
+ uint hash_mode = HASH_MODE;
+ uint attack_mode = ATTACK_MODE;
+ uint markov_disable = MARKOV_DISABLE;
+ uint markov_classic = MARKOV_CLASSIC;
+ uint markov_threshold = MARKOV_THRESHOLD;
+ char *markov_hcstat = NULL;
+ char *outfile = NULL;
+ uint outfile_format = OUTFILE_FORMAT;
+ uint outfile_autohex = OUTFILE_AUTOHEX;
+ uint outfile_check_timer = OUTFILE_CHECK_TIMER;
+ uint restore = RESTORE;
+ uint restore_timer = RESTORE_TIMER;
+ uint restore_disable = RESTORE_DISABLE;
+ uint status = STATUS;
+ uint status_timer = STATUS_TIMER;
+ uint machine_readable = MACHINE_READABLE;
+ uint loopback = LOOPBACK;
+ uint weak_hash_threshold = WEAK_HASH_THRESHOLD;
+ char *session = NULL;
+ uint hex_charset = HEX_CHARSET;
+ uint hex_salt = HEX_SALT;
+ uint hex_wordlist = HEX_WORDLIST;
+ uint rp_gen = RP_GEN;
+ uint rp_gen_func_min = RP_GEN_FUNC_MIN;
+ uint rp_gen_func_max = RP_GEN_FUNC_MAX;
+ uint rp_gen_seed = RP_GEN_SEED;
+ char *rule_buf_l = (char *) RULE_BUF_L;
+ char *rule_buf_r = (char *) RULE_BUF_R;
+ uint increment = INCREMENT;
+ uint increment_min = INCREMENT_MIN;
+ uint increment_max = INCREMENT_MAX;
+ char *cpu_affinity = NULL;
+ OCL_PTR *ocl = NULL;
+ char *opencl_devices = NULL;
+ char *opencl_platforms = NULL;
+ 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;
+ uint gpu_temp_disable = GPU_TEMP_DISABLE;
#ifdef HAVE_HWMON
- uint gpu_temp_abort = GPU_TEMP_ABORT;
- uint gpu_temp_retain = GPU_TEMP_RETAIN;
+ uint gpu_temp_abort = GPU_TEMP_ABORT;
+ uint gpu_temp_retain = GPU_TEMP_RETAIN;
#ifdef HAVE_ADL
- uint powertune_enable = POWERTUNE_ENABLE;
+ uint powertune_enable = POWERTUNE_ENABLE;
#endif
#endif
- uint logfile_disable = LOGFILE_DISABLE;
- uint segment_size = SEGMENT_SIZE;
- uint scrypt_tmto = SCRYPT_TMTO;
- char separator = SEPARATOR;
- uint bitmap_min = BITMAP_MIN;
- uint bitmap_max = BITMAP_MAX;
- char *custom_charset_1 = NULL;
- char *custom_charset_2 = NULL;
- char *custom_charset_3 = NULL;
- char *custom_charset_4 = NULL;
+ uint logfile_disable = LOGFILE_DISABLE;
+ uint segment_size = SEGMENT_SIZE;
+ uint scrypt_tmto = SCRYPT_TMTO;
+ char separator = SEPARATOR;
+ uint bitmap_min = BITMAP_MIN;
+ uint bitmap_max = BITMAP_MAX;
+ char *custom_charset_1 = NULL;
+ char *custom_charset_2 = NULL;
+ char *custom_charset_3 = NULL;
+ char *custom_charset_4 = NULL;
#define IDX_HELP 'h'
#define IDX_VERSION 'V'
#define IDX_RESTORE_DISABLE 0xff27
#define IDX_STATUS 0xff17
#define IDX_STATUS_TIMER 0xff18
- #define IDX_STATUS_AUTOMAT 0xff50
+ #define IDX_MACHINE_READABLE 0xff50
#define IDX_LOOPBACK 0xff38
#define IDX_WEAK_HASH_THRESHOLD 0xff42
#define IDX_SESSION 0xff19
{"restore-disable", no_argument, 0, IDX_RESTORE_DISABLE},
{"status", no_argument, 0, IDX_STATUS},
{"status-timer", required_argument, 0, IDX_STATUS_TIMER},
- {"status-automat", no_argument, 0, IDX_STATUS_AUTOMAT},
+ {"machine-readable", no_argument, 0, IDX_MACHINE_READABLE},
{"loopback", no_argument, 0, IDX_LOOPBACK},
{"weak-hash-threshold", required_argument, 0, IDX_WEAK_HASH_THRESHOLD},
{"session", required_argument, 0, IDX_SESSION},
case IDX_RESTORE_DISABLE: restore_disable = 1; break;
case IDX_STATUS: status = 1; break;
case IDX_STATUS_TIMER: status_timer = atoi (optarg); break;
- case IDX_STATUS_AUTOMAT: status_automat = 1; break;
+ case IDX_MACHINE_READABLE: machine_readable = 1; break;
case IDX_LOOPBACK: loopback = 1; break;
case IDX_WEAK_HASH_THRESHOLD: weak_hash_threshold = atoi (optarg); break;
//case IDX_SESSION: session = optarg; break;
{
if (benchmark == 1)
{
- if (status_automat == 0)
+ if (machine_readable == 0)
{
log_info ("%s (%s) starting in benchmark-mode...", PROGNAME, VERSION_TAG);
log_info ("");
}
}
- if ((workload_profile < 1) || (workload_profile > 3))
+ if ((workload_profile < 1) || (workload_profile > 4))
{
log_error ("ERROR: workload-profile %i not available", workload_profile);
* store stuff
*/
- data.hash_mode = hash_mode;
- data.restore = restore;
- data.restore_timer = restore_timer;
- data.restore_disable = restore_disable;
- data.status = status;
- data.status_timer = status_timer;
- data.status_automat = status_automat;
- data.loopback = loopback;
- data.runtime = runtime;
- data.remove = remove;
- data.remove_timer = remove_timer;
- data.debug_mode = debug_mode;
- data.debug_file = debug_file;
- data.username = username;
- data.quiet = quiet;
- data.outfile = outfile;
- data.outfile_format = outfile_format;
- data.outfile_autohex = outfile_autohex;
- data.hex_charset = hex_charset;
- data.hex_salt = hex_salt;
- data.hex_wordlist = hex_wordlist;
- data.separator = separator;
- data.rp_files = rp_files;
- data.rp_files_cnt = rp_files_cnt;
- data.rp_gen = rp_gen;
- data.rp_gen_seed = rp_gen_seed;
- data.force = force;
- data.benchmark = benchmark;
- data.skip = skip;
- data.limit = limit;
+ data.hash_mode = hash_mode;
+ data.restore = restore;
+ data.restore_timer = restore_timer;
+ data.restore_disable = restore_disable;
+ data.status = status;
+ data.status_timer = status_timer;
+ data.machine_readable = machine_readable;
+ data.loopback = loopback;
+ data.runtime = runtime;
+ data.remove = remove;
+ data.remove_timer = remove_timer;
+ data.debug_mode = debug_mode;
+ data.debug_file = debug_file;
+ data.username = username;
+ data.quiet = quiet;
+ data.outfile = outfile;
+ data.outfile_format = outfile_format;
+ data.outfile_autohex = outfile_autohex;
+ data.hex_charset = hex_charset;
+ data.hex_salt = hex_salt;
+ data.hex_wordlist = hex_wordlist;
+ data.separator = separator;
+ data.rp_files = rp_files;
+ data.rp_files_cnt = rp_files_cnt;
+ data.rp_gen = rp_gen;
+ data.rp_gen_seed = rp_gen_seed;
+ data.force = force;
+ data.benchmark = benchmark;
+ data.skip = skip;
+ data.limit = limit;
#ifdef HAVE_HWMON
#ifdef HAVE_ADL
- data.powertune_enable = powertune_enable;
+ data.powertune_enable = powertune_enable;
#endif
#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;
+ 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;
/**
* cpu affinity
logfile_top_uint (segment_size);
logfile_top_uint (show);
logfile_top_uint (status);
- logfile_top_uint (status_automat);
+ logfile_top_uint (machine_readable);
logfile_top_uint (status_timer);
logfile_top_uint (usage);
logfile_top_uint (username);
weak_hash_threshold = 0;
gpu_temp_disable = 1;
+ #ifdef HAVE_HWMON
+ #ifdef HAVE_ADL
+ powertune_enable = 1;
+ #endif
+ #endif
+
data.status_timer = status_timer;
data.restore_timer = restore_timer;
data.restore_disable = restore_disable;
break;
case 7400: if (pw_max > 16) pw_max = 16;
break;
+ case 7500: if (pw_max > 8) pw_max = 8;
+ break;
case 7900: if (pw_max > 48) pw_max = 48;
break;
case 8500: if (pw_max > 8) pw_max = 8;
// this causes trouble with vendor id based macros
// we'll assign generic to those without special optimization available
- cl_uint vendor_id = 0;
+ cl_uint platform_vendor_id = 0;
if (strcmp (platform_vendor, CL_VENDOR_AMD) == 0)
{
- vendor_id = VENDOR_ID_AMD;
+ platform_vendor_id = VENDOR_ID_AMD;
+ }
+ else if (strcmp (platform_vendor, CL_VENDOR_AMD_USE_INTEL) == 0)
+ {
+ platform_vendor_id = VENDOR_ID_AMD_USE_INTEL;
}
else if (strcmp (platform_vendor, CL_VENDOR_APPLE) == 0)
{
- vendor_id = VENDOR_ID_APPLE;
+ platform_vendor_id = VENDOR_ID_APPLE;
}
else if (strcmp (platform_vendor, CL_VENDOR_INTEL_BEIGNET) == 0)
{
- vendor_id = VENDOR_ID_INTEL_BEIGNET;
+ platform_vendor_id = VENDOR_ID_INTEL_BEIGNET;
}
else if (strcmp (platform_vendor, CL_VENDOR_INTEL_SDK) == 0)
{
- vendor_id = VENDOR_ID_INTEL_SDK;
+ platform_vendor_id = VENDOR_ID_INTEL_SDK;
}
else if (strcmp (platform_vendor, CL_VENDOR_MESA) == 0)
{
- vendor_id = VENDOR_ID_MESA;
+ platform_vendor_id = VENDOR_ID_MESA;
}
else if (strcmp (platform_vendor, CL_VENDOR_NV) == 0)
{
- vendor_id = VENDOR_ID_NV;
+ platform_vendor_id = VENDOR_ID_NV;
}
else if (strcmp (platform_vendor, CL_VENDOR_POCL) == 0)
{
- vendor_id = VENDOR_ID_POCL;
+ platform_vendor_id = VENDOR_ID_POCL;
}
else
{
- vendor_id = VENDOR_ID_GENERIC;
+ platform_vendor_id = VENDOR_ID_GENERIC;
}
for (uint platform_devices_id = 0; platform_devices_id < platform_devices_cnt; platform_devices_id++)
hc_device_param_t *device_param = &data.devices_param[device_id];
- device_param->vendor_id = vendor_id;
+ device_param->platform_vendor_id = platform_vendor_id;
device_param->device = platform_devices[platform_devices_id];
device_param->device_name = device_name;
+ // device_vendor
+
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR, 0, NULL, ¶m_value_size);
+
+ char *device_vendor = (char *) mymalloc (param_value_size);
+
+ hc_clGetDeviceInfo (data.ocl, device_param->device, CL_DEVICE_VENDOR, param_value_size, device_vendor, NULL);
+
+ device_param->device_vendor = device_vendor;
+
+ cl_uint device_vendor_id = 0;
+
+ if (strcmp (device_vendor, CL_VENDOR_AMD) == 0)
+ {
+ device_vendor_id = VENDOR_ID_AMD;
+ }
+ else if (strcmp (device_vendor, CL_VENDOR_AMD_USE_INTEL) == 0)
+ {
+ device_vendor_id = VENDOR_ID_AMD_USE_INTEL;
+ }
+ else if (strcmp (device_vendor, CL_VENDOR_APPLE) == 0)
+ {
+ device_vendor_id = VENDOR_ID_APPLE;
+ }
+ else if (strcmp (device_vendor, CL_VENDOR_INTEL_BEIGNET) == 0)
+ {
+ device_vendor_id = VENDOR_ID_INTEL_BEIGNET;
+ }
+ else if (strcmp (device_vendor, CL_VENDOR_INTEL_SDK) == 0)
+ {
+ device_vendor_id = VENDOR_ID_INTEL_SDK;
+ }
+ else if (strcmp (device_vendor, CL_VENDOR_MESA) == 0)
+ {
+ device_vendor_id = VENDOR_ID_MESA;
+ }
+ else if (strcmp (device_vendor, CL_VENDOR_NV) == 0)
+ {
+ device_vendor_id = VENDOR_ID_NV;
+ }
+ else if (strcmp (device_vendor, CL_VENDOR_POCL) == 0)
+ {
+ device_vendor_id = VENDOR_ID_POCL;
+ }
+ else
+ {
+ device_vendor_id = VENDOR_ID_GENERIC;
+ }
+
+ device_param->device_vendor_id = device_vendor_id;
+
// tuning db
tuning_db_entry_t *tuningdb_entry = tuning_db_search (tuning_db, device_param, attack_mode, hash_mode);
if (device_endian_little == CL_FALSE)
{
- log_info ("Device #%u: WARNING: not little endian device", device_id + 1);
+ if (data.quiet == 0) log_info ("Device #%u: WARNING: not little endian device", device_id + 1);
device_param->skipped = 1;
}
if (device_available == CL_FALSE)
{
- log_info ("Device #%u: WARNING: device not available", device_id + 1);
+ if (data.quiet == 0) log_info ("Device #%u: WARNING: device not available", device_id + 1);
device_param->skipped = 1;
}
if (device_compiler_available == CL_FALSE)
{
- log_info ("Device #%u: WARNING: device no compiler available", device_id + 1);
+ if (data.quiet == 0) log_info ("Device #%u: WARNING: device no compiler available", device_id + 1);
device_param->skipped = 1;
}
if ((device_execution_capabilities & CL_EXEC_KERNEL) == 0)
{
- log_info ("Device #%u: WARNING: device does not support executing kernels", device_id + 1);
+ if (data.quiet == 0) log_info ("Device #%u: WARNING: device does not support executing kernels", device_id + 1);
device_param->skipped = 1;
}
if (strstr (device_extensions, "base_atomics") == 0)
{
- log_info ("Device #%u: WARNING: device does not support base atomics", device_id + 1);
+ if (data.quiet == 0) log_info ("Device #%u: WARNING: device does not support base atomics", device_id + 1);
device_param->skipped = 1;
}
if (strstr (device_extensions, "byte_addressable_store") == 0)
{
- log_info ("Device #%u: WARNING: device does not support byte addressable store", device_id + 1);
+ if (data.quiet == 0) log_info ("Device #%u: WARNING: device does not support byte addressable store", device_id + 1);
device_param->skipped = 1;
}
if (device_local_mem_size < 32768)
{
- log_info ("Device #%u: WARNING: device local mem size is too small", device_id + 1);
+ if (data.quiet == 0) log_info ("Device #%u: WARNING: device local mem size is too small", device_id + 1);
device_param->skipped = 1;
}
+ // If there's both an Intel CPU and an AMD OpenCL runtime it's a tricky situation
+ // Both platforms support CPU device types and therefore both will try to use 100% of the physical resources
+ // This results in both utilizing it for 50%
+ // However, Intel has much better SIMD control over their own hardware
+ // It makes sense to give them full control over their own hardware
+
+ if (device_type & CL_DEVICE_TYPE_CPU)
+ {
+ if (device_param->device_vendor_id == VENDOR_ID_AMD_USE_INTEL)
+ {
+ if (data.quiet == 0) log_info ("Device #%u: WARNING: not native intel opencl platform", device_id + 1);
+
+ device_param->skipped = 1;
+ }
+ }
// skipped
char *device_name_chksum = (char *) mymalloc (INFOSZ);
#if __x86_64__
- snprintf (device_name_chksum, INFOSZ - 1, "%u-%u-%u-%s-%s-%s-%u", 64, device_param->vendor_id, device_param->vector_width, device_param->device_name, device_param->device_version, device_param->driver_version, COMPTIME);
+ snprintf (device_name_chksum, INFOSZ - 1, "%u-%u-%u-%s-%s-%s-%u", 64, device_param->platform_vendor_id, device_param->vector_width, device_param->device_name, device_param->device_version, device_param->driver_version, COMPTIME);
#else
- snprintf (device_name_chksum, INFOSZ - 1, "%u-%u-%u-%s-%s-%s-%u", 32, device_param->vendor_id, device_param->vector_width, device_param->device_name, device_param->device_version, device_param->driver_version, COMPTIME);
+ snprintf (device_name_chksum, INFOSZ - 1, "%u-%u-%u-%s-%s-%s-%u", 32, device_param->platform_vendor_id, device_param->vector_width, device_param->device_name, device_param->device_version, device_param->driver_version, COMPTIME);
#endif
uint device_name_digest[4] = { 0 };
if (device_type & CL_DEVICE_TYPE_GPU)
{
- if (vendor_id == VENDOR_ID_AMD)
+ if (device_vendor_id == VENDOR_ID_AMD)
{
cl_uint device_processor_cores = 0;
device_param->device_processor_cores = device_processor_cores;
}
- else if (vendor_id == VENDOR_ID_NV)
+ else if (device_vendor_id == VENDOR_ID_NV)
{
cl_uint kernel_exec_timeout = 0;
if ((benchmark == 1 || quiet == 0) && (algorithm_pos == 0))
{
- if (status_automat == 0)
+ if (machine_readable == 0)
{
if (device_param->skipped == 0)
{
- log_info ("Device #%u: %s, %lu/%lu MB allocatable, %dMhz, %uMCU",
+ log_info ("Device #%u: %s, %lu/%lu MB allocatable, %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
{
if (device_type & CL_DEVICE_TYPE_GPU)
{
- if (vendor_id == VENDOR_ID_AMD)
+ if (platform_vendor_id == VENDOR_ID_AMD)
{
int catalyst_check = (force == 1) ? 0 : 1;
return (-1);
}
}
- else if (vendor_id == VENDOR_ID_NV)
+ else if (platform_vendor_id == VENDOR_ID_NV)
{
if (device_param->kernel_exec_timeout != 0)
{
}
}
+ /* turns out pocl still creates segfaults (because of llvm)
if (device_type & CL_DEVICE_TYPE_CPU)
{
- if (vendor_id == VENDOR_ID_AMD)
+ if (platform_vendor_id == VENDOR_ID_AMD)
{
if (force == 0)
{
}
}
}
+ */
/**
* kernel accel and loops tuning db adjustment
if ((benchmark == 1 || quiet == 0) && (algorithm_pos == 0))
{
- if (status_automat == 0)
+ if (machine_readable == 0)
{
log_info ("");
}
#ifdef HAVE_HWMON
#if defined(HAVE_NVML) || defined(HAVE_NVAPI)
- hm_attrs_t hm_adapters_nv[DEVICES_MAX] = { { { 0 }, 0, 0 } };
+ hm_attrs_t hm_adapters_nv[DEVICES_MAX] = { { { 0 }, 0, 0, 0, 0 } };
#endif
#ifdef HAVE_ADL
- hm_attrs_t hm_adapters_amd[DEVICES_MAX] = { { { 0 }, 0, 0 } };
+ hm_attrs_t hm_adapters_amd[DEVICES_MAX] = { { { 0 }, 0, 0, 0, 0 } };
#endif
if (gpu_temp_disable == 0)
const uint platform_devices_id = device_param->platform_devices_id;
#if defined(HAVE_NVML) || defined(HAVE_NVAPI)
- if (device_param->vendor_id == VENDOR_ID_NV)
+ if (device_param->device_vendor_id == VENDOR_ID_NV)
{
memcpy (&data.hm_device[device_id], &hm_adapters_nv[platform_devices_id], sizeof (hm_attrs_t));
}
#endif
#ifdef HAVE_ADL
- if (device_param->vendor_id == VENDOR_ID_AMD)
+ if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
memcpy (&data.hm_device[device_id], &hm_adapters_amd[platform_devices_id], sizeof (hm_attrs_t));
}
}
}
- /*
- * Temporary fix:
- * with AMD r9 295x cards it seems that we need to set the powertune value just AFTER the ocl init stuff
- * otherwise after hc_clCreateContext () etc, powertune value was set back to "normal" and cards unfortunately
- * were not working @ full speed (setting hm_ADL_Overdrive_PowerControl_Set () here seems to fix the problem)
- * Driver / ADL bug?
- */
+ /**
+ * Temporary fix:
+ * with AMD r9 295x cards it seems that we need to set the powertune value just AFTER the ocl init stuff
+ * otherwise after hc_clCreateContext () etc, powertune value was set back to "normal" and cards unfortunately
+ * were not working @ full speed (setting hm_ADL_Overdrive_PowerControl_Set () here seems to fix the problem)
+ * Driver / ADL bug?
+ */
#ifdef HAVE_ADL
if (powertune_enable == 1)
if (hash_mode == 8900)
{
- if (device_param->vendor_id == VENDOR_ID_AMD)
+ if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
tmto_start = 1;
}
- else if (device_param->vendor_id == VENDOR_ID_NV)
+ else if (device_param->device_vendor_id == VENDOR_ID_NV)
{
tmto_start = 2;
}
}
else if (hash_mode == 9300)
{
- if (device_param->vendor_id == VENDOR_ID_AMD)
+ if (device_param->device_vendor_id == VENDOR_ID_AMD)
{
tmto_start = 2;
}
- else if (device_param->vendor_id == VENDOR_ID_NV)
+ else if (device_param->device_vendor_id == VENDOR_ID_NV)
{
tmto_start = 2;
}
int skip = 0;
- if (size_pws > device_param->device_maxmem_alloc) skip = 1;
- if (size_tmps > device_param->device_maxmem_alloc) skip = 1;
- if (size_hooks > device_param->device_maxmem_alloc) skip = 1;
-
- if (( bitmap_size
- + bitmap_size
- + bitmap_size
- + bitmap_size
- + bitmap_size
- + bitmap_size
- + bitmap_size
- + bitmap_size
- + size_bfs
- + size_combs
- + size_digests
- + size_esalts
- + size_hooks
- + size_markov_css
- + size_plains
- + size_pws
- + size_pws // not a bug
- + size_results
- + size_root_css
- + size_rules
- + size_rules_c
- + size_salts
- + size_scryptV
- + size_shown
- + size_tm
- + size_tmps) > device_param->device_global_mem) skip = 1;
+ const u64 size_total
+ = bitmap_size
+ + bitmap_size
+ + bitmap_size
+ + bitmap_size
+ + bitmap_size
+ + bitmap_size
+ + bitmap_size
+ + bitmap_size
+ + size_bfs
+ + size_combs
+ + size_digests
+ + size_esalts
+ + size_hooks
+ + size_markov_css
+ + size_plains
+ + size_pws
+ + size_pws // not a bug
+ + size_results
+ + size_root_css
+ + size_rules
+ + size_rules_c
+ + size_salts
+ + size_scryptV
+ + size_shown
+ + size_tm
+ + size_tmps;
+
+ // Don't ask me, ask AMD!
+
+ if (size_total > device_param->device_maxmem_alloc) skip = 1;
+ if (size_total > device_param->device_global_mem) skip = 1;
if (skip == 1)
{
// we don't have sm_* on vendors not NV but it doesn't matter
- 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 _WIN
+ snprintf (build_opts, sizeof (build_opts) - 1, "-I \"%s\\OpenCL\\\" -I '%s\\OpenCL\\' -I %s\\OpenCL\\ -I\"%s\\OpenCL\\\" -I'%s\\OpenCL\\' -I%s\\OpenCL\\", shared_dir, shared_dir, shared_dir, shared_dir, shared_dir, shared_dir);
+ #else
+ snprintf (build_opts, sizeof (build_opts) - 1, "-I \"%s/OpenCL/\" -I '%s/OpenCL/' -I %s/OpenCL/ -I\"%s/OpenCL/\" -I'%s/OpenCL/' -I%s/OpenCL/", shared_dir, shared_dir, shared_dir, shared_dir, shared_dir, shared_dir);
+ #endif
+
+ char build_opts_new[1024] = { 0 };
+
+ snprintf (build_opts_new, sizeof (build_opts_new) - 1, "%s -DVENDOR_ID=%u -DCUDA_ARCH=%d -DVECT_SIZE=%u -DDEVICE_TYPE=%u -DKERN_TYPE=%u -D_unroll -cl-std=CL1.1", build_opts, device_param->device_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)
+ strncpy (build_opts, build_opts_new, sizeof (build_opts) - 1);
+
+ /*
+ if (device_param->device_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);
hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_digests_shown, CL_TRUE, 0, size_shown, data.digests_shown, 0, NULL, NULL);
hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_salt_bufs, CL_TRUE, 0, size_salts, data.salts_buf, 0, NULL, NULL);
- run_kernel_bzero (device_param, device_param->d_pws_buf, size_pws);
- run_kernel_bzero (device_param, device_param->d_pws_amp_buf, size_pws);
- run_kernel_bzero (device_param, device_param->d_tmps, size_tmps);
- run_kernel_bzero (device_param, device_param->d_hooks, size_hooks);
- run_kernel_bzero (device_param, device_param->d_plain_bufs, size_plains);
- run_kernel_bzero (device_param, device_param->d_result, size_results);
-
/**
* special buffers
*/
device_param->d_rules_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_rules_c, NULL);
hc_clEnqueueWriteBuffer (data.ocl, device_param->command_queue, device_param->d_rules, CL_TRUE, 0, size_rules, kernel_rules_buf, 0, NULL, NULL);
-
- run_kernel_bzero (device_param, device_param->d_rules_c, size_rules_c);
}
else if (attack_kern == ATTACK_KERN_COMBI)
{
device_param->d_combs_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_combs, NULL);
device_param->d_root_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
device_param->d_markov_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
-
- run_kernel_bzero (device_param, device_param->d_combs, size_combs);
- run_kernel_bzero (device_param, device_param->d_combs_c, size_combs);
- run_kernel_bzero (device_param, device_param->d_root_css_buf, size_root_css);
- run_kernel_bzero (device_param, device_param->d_markov_css_buf, size_markov_css);
}
else if (attack_kern == ATTACK_KERN_BF)
{
device_param->d_tm_c = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_tm, NULL);
device_param->d_root_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_root_css, NULL);
device_param->d_markov_css_buf = hc_clCreateBuffer (data.ocl, device_param->context, CL_MEM_READ_ONLY, size_markov_css, NULL);
-
- run_kernel_bzero (device_param, device_param->d_bfs, size_bfs);
- run_kernel_bzero (device_param, device_param->d_bfs_c, size_bfs);
- run_kernel_bzero (device_param, device_param->d_tm_c, size_tm);
- run_kernel_bzero (device_param, device_param->d_root_css_buf, size_root_css);
- run_kernel_bzero (device_param, device_param->d_markov_css_buf, size_markov_css);
}
if (size_esalts)
device_param->kernel_params_tm[0] = &device_param->d_bfs_c;
device_param->kernel_params_tm[1] = &device_param->d_tm_c;
+ device_param->kernel_params_memset_buf32[1] = 0; // value
+ device_param->kernel_params_memset_buf32[2] = 0; // gid_max
+
+ device_param->kernel_params_memset[0] = NULL;
+ device_param->kernel_params_memset[1] = &device_param->kernel_params_memset_buf32[1];
+ device_param->kernel_params_memset[2] = &device_param->kernel_params_memset_buf32[2];
+
/**
* kernel name
*/
if (opts_type & OPTS_TYPE_HOOK23) hc_clSetKernelArg (data.ocl, device_param->kernel23, i, sizeof (cl_uint), device_param->kernel_params[i]);
}
+ // GPU memset
+
+ device_param->kernel_memset = hc_clCreateKernel (data.ocl, device_param->program, "gpu_memset");
+
+ hc_clGetKernelWorkGroupInfo (data.ocl, device_param->kernel_memset, device_param->device, CL_KERNEL_WORK_GROUP_SIZE, sizeof (size_t), &kernel_wgs_tmp, NULL); kernel_threads = MIN (kernel_threads, kernel_wgs_tmp);
+
+ hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 0, sizeof (cl_mem), device_param->kernel_params_memset[0]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 1, sizeof (cl_uint), device_param->kernel_params_memset[1]);
+ hc_clSetKernelArg (data.ocl, device_param->kernel_memset, 2, sizeof (cl_uint), device_param->kernel_params_memset[2]);
+
+ // MP start
+
if (attack_mode == ATTACK_MODE_BF)
{
device_param->kernel_mp_l = hc_clCreateKernel (data.ocl, device_param->program_mp, "l_markov");
device_param->kernel_threads = kernel_threads;
+ // zero some data buffers
+
+ run_kernel_bzero (device_param, device_param->d_pws_buf, size_pws);
+ run_kernel_bzero (device_param, device_param->d_pws_amp_buf, size_pws);
+ run_kernel_bzero (device_param, device_param->d_tmps, size_tmps);
+ run_kernel_bzero (device_param, device_param->d_hooks, size_hooks);
+ run_kernel_bzero (device_param, device_param->d_plain_bufs, size_plains);
+ run_kernel_bzero (device_param, device_param->d_result, size_results);
+
+ /**
+ * special buffers
+ */
+
+ if (attack_kern == ATTACK_KERN_STRAIGHT)
+ {
+ run_kernel_bzero (device_param, device_param->d_rules_c, size_rules_c);
+ }
+ else if (attack_kern == ATTACK_KERN_COMBI)
+ {
+ run_kernel_bzero (device_param, device_param->d_combs, size_combs);
+ run_kernel_bzero (device_param, device_param->d_combs_c, size_combs);
+ run_kernel_bzero (device_param, device_param->d_root_css_buf, size_root_css);
+ run_kernel_bzero (device_param, device_param->d_markov_css_buf, size_markov_css);
+ }
+ else if (attack_kern == ATTACK_KERN_BF)
+ {
+ run_kernel_bzero (device_param, device_param->d_bfs, size_bfs);
+ run_kernel_bzero (device_param, device_param->d_bfs_c, size_bfs);
+ run_kernel_bzero (device_param, device_param->d_tm_c, size_tm);
+ run_kernel_bzero (device_param, device_param->d_root_css_buf, size_root_css);
+ run_kernel_bzero (device_param, device_param->d_markov_css_buf, size_markov_css);
+ }
+
+ /**
+ * Store thermal target temperature so we can send a notice to user
+ */
+
+ #if defined(HAVE_HWMON)
+ if (gpu_temp_disable == 0)
+ {
+ const int gpu_temp_threshold_slowdown = hm_get_threshold_slowdown_with_device_id (device_id);
+
+ data.hm_device[device_id].gpu_temp_threshold_slowdown = (gpu_temp_threshold_slowdown == -1) ? 100000 : gpu_temp_threshold_slowdown;
+ }
+ #endif
+
/**
* Store initial fanspeed if gpu_temp_retain is enabled
*/
if (benchmark == 1)
{
- if (status_automat == 0)
+ if (machine_readable == 0)
{
quiet = 0;
{
status_benchmark ();
- if (status_automat == 0)
+ if (machine_readable == 0)
{
log_info ("");
}
if (device_param->kernel_mp_r) hc_clReleaseKernel (data.ocl, device_param->kernel_mp_r);
if (device_param->kernel_tm) hc_clReleaseKernel (data.ocl, device_param->kernel_tm);
if (device_param->kernel_amp) hc_clReleaseKernel (data.ocl, device_param->kernel_amp);
+ if (device_param->kernel_memset) hc_clReleaseKernel (data.ocl, device_param->kernel_memset);
if (device_param->program) hc_clReleaseProgram (data.ocl, device_param->program);
if (device_param->program_mp) hc_clReleaseProgram (data.ocl, device_param->program_mp);
projects / hashcat.git / blobdiff