const uint VERSION_BIN = 300;
const uint RESTORE_MIN = 300;
-double TARGET_MS_PROFILE[3] = { 8, 16, 96 };
+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 STATUS_AUTOMATE 0
#define LOOPBACK 0
#define WEAK_HASH_THRESHOLD 100
#define SHOW 0
#define MAX_DICTSTAT 10000
-#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 137
+#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 143
#define global_free(attr) \
{ \
6221,
6231,
6241,
+ 13711,
+ 13721,
+ 13731,
+ 13741,
+ 13751,
+ 13761,
8800,
12900,
12200,
"",
"Usage: %s [options]... hash|hashfile|hccapfile [dictionary|mask|directory]...",
"",
- "=======",
- "Options",
- "=======",
+ "- [ Options ] -",
"",
- "* General:",
- "",
- " -m, --hash-type=NUM Hash-type, see references below",
- " -a, --attack-mode=NUM Attack-mode, see references below",
- " -V, --version Print version",
- " -h, --help Print help",
- " --quiet Suppress output",
- "",
- "* Misc:",
- "",
- " --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 Seconds between status-screen update",
- " --status-automat Display the status view in a machine readable format",
- " --loopback Add new plains to induct directory",
- " --weak-hash-threshold=NUM Threshold when to stop checking for weak hashes, default is 100 salts",
- "",
- "* Markov:",
- "",
- " --markov-hcstat=FILE Specify hcstat file to use, default is hashcat.hcstat",
- " --markov-disable Disables markov-chains, emulates classic brute-force",
- " --markov-classic Enables classic markov-chains, no per-position enhancement",
- " -t, --markov-threshold=NUM Threshold when to stop accepting new markov-chains",
- "",
- "* Session:",
- "",
- " --runtime=NUM Abort session after NUM seconds of runtime",
- " --session=STR Define specific session name",
- " --restore Restore session from --session",
- " --restore-disable Do not write restore file",
- "",
- "* Files:",
- "",
- " -o, --outfile=FILE Define outfile for recovered hash",
- " --outfile-format=NUM Define outfile-format for recovered hash, see references below",
- " --outfile-autohex-disable Disable the use of $HEX[] in output plains",
- " --outfile-check-timer=NUM Seconds between outfile checks",
- " -p, --separator=CHAR Separator char for hashlists and outfile",
- " --show Show cracked passwords only",
- " --left Show un-cracked passwords only",
- " --username Enable ignoring of usernames in hashfile (recommended: also use --show)",
- " --remove Enable remove of hash once it is cracked",
- " --remove-timer=NUM Update input hash file each NUM seconds",
- " --potfile-disable Do not write potfile",
- " --potfile-path Specific path to potfile",
- " --debug-mode=NUM Defines the debug mode (hybrid only by using rules), see references below",
- " --debug-file=FILE Output file for debugging rules (see also --debug-mode)",
- " --induction-dir=FOLDER Specify induction directory to use, default is $session.induct",
- " --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",
- "",
- "* Resources:",
- "",
- " -b, --benchmark Run benchmark",
- " --benchmark-repeats=NUM Repeat the kernel on the device NUM times to increase benchmark accuracy",
- " -c, --segment-size=NUM Size in MB to cache from the wordfile",
- " --bitmap-min=NUM Minimum number of bits allowed for bitmaps",
- " --bitmap-max=NUM Maximum number of bits allowed for bitmaps",
- " --cpu-affinity=STR Locks to CPU devices, separate with comma",
- " --opencl-platforms=STR OpenCL platforms to use, separate with comma",
- " -d, --opencl-devices=STR OpenCL devices to use, separate with comma",
- " --opencl-device-types=STR OpenCL device-types to use, separate with comma, see references below",
- " --opencl-vector-width=NUM OpenCL vector-width (either 1, 2, 4, 8 or 16), overrides value from device query",
- " -w, --workload-profile=NUM Enable a specific workload profile, see references below",
- " -n, --kernel-accel=NUM Workload tuning, increase the outer-loop step size",
- " -u, --kernel-loops=NUM Workload tuning, increase the inner-loop step size",
- " --gpu-temp-disable Disable temperature and fanspeed readings and triggers",
+ " 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",
+ " --status-automate | | 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=NUM Abort session if GPU temperature reaches NUM degrees celsius",
- " --gpu-temp-retain=NUM Try to retain GPU temperature at NUM degrees celsius (AMD only)",
+ " --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 option (AMD OverDrive 6 only)",
+ " --powertune-enable | | Enable automatic power tuning (AMD OverDrive 6 only) |",
#endif
#endif
- " --scrypt-tmto=NUM Manually override automatically calculated TMTO value for scrypt",
- "",
- "* Distributed:",
- "",
- " -s, --skip=NUM Skip number of words",
- " -l, --limit=NUM Limit number of words",
- " --keyspace Show keyspace base:mod values and quit",
- "",
- "* Rules:",
- "",
- " -j, --rule-left=RULE Single rule applied to each word from left dict",
- " -k, --rule-right=RULE Single rule applied to each word from right dict",
- " -r, --rules-file=FILE Rules-file, multi use: -r 1.rule -r 2.rule",
- " -g, --generate-rules=NUM Generate NUM random rules",
- " --generate-rules-func-min=NUM Force NUM functions per random rule min",
- " --generate-rules-func-max=NUM Force NUM functions per random rule max",
- " --generate-rules-seed=NUM Force RNG seed to NUM",
- "",
- "* Custom charsets:",
- "",
- " -1, --custom-charset1=CS User-defined charsets",
- " -2, --custom-charset2=CS Example:",
- " -3, --custom-charset3=CS --custom-charset1=?dabcdef : sets charset ?1 to 0123456789abcdef",
- " -4, --custom-charset4=CS -2 mycharset.hcchr : sets charset ?2 to chars contained in file",
- "",
- "* Increment:",
- "",
- " -i, --increment Enable increment mode",
- " --increment-min=NUM Start incrementing at NUM",
- " --increment-max=NUM Stop incrementing at NUM",
- "",
- "==========",
- "References",
- "==========",
- "",
- "* Workload Profile:",
- "",
- " 1 = Interactive performance profile, kernel execution runtime to 8ms, lower latency desktop, lower speed",
- " 2 = Default performance profile, kernel execution runtime to 16ms, economic setting",
- " 3 = Headless performance profile, kernel execution runtime to 96ms, higher latency desktop, higher speed",
- "",
- "* OpenCL device-types:",
- "",
- " 1 = CPU devices",
- " 2 = GPU devices",
- " 3 = Accelerator devices (FPGA, CELL Blade, etc.)",
- "",
- "* Outfile Formats:",
- "",
- " 1 = hash[:salt]",
- " 2 = plain",
- " 3 = hash[:salt]:plain",
- " 4 = hex_plain",
- " 5 = hash[:salt]:hex_plain",
- " 6 = plain:hex_plain",
- " 7 = hash[:salt]:plain:hex_plain",
- " 8 = crackpos",
- " 9 = hash[:salt]:crackpos",
- " 10 = plain:crackpos",
- " 11 = hash[:salt]:plain:crackpos",
- " 12 = hex_plain:crackpos",
- " 13 = hash[:salt]:hex_plain:crackpos",
- " 14 = plain:hex_plain:crackpos",
- " 15 = hash[:salt]:plain:hex_plain:crackpos",
- "",
- "* Debug mode output formats (for hybrid mode only, by using rules):",
- "",
- " 1 = save finding rule",
- " 2 = save original word",
- " 3 = save original word and finding rule",
- " 4 = save original word, finding rule and modified plain",
- "",
- "* Built-in charsets:",
- "",
- " ?l = abcdefghijklmnopqrstuvwxyz",
- " ?u = ABCDEFGHIJKLMNOPQRSTUVWXYZ",
- " ?d = 0123456789",
- " ?s = !\"#$%%&'()*+,-./:;<=>?@[\\]^_`{|}~",
- " ?a = ?l?u?d?s",
- " ?b = 0x00 - 0xff",
- "",
- "* Attack modes:",
- "",
- " 0 = Straight",
- " 1 = Combination",
- " 3 = Brute-force",
- " 6 = Hybrid dict + mask",
- " 7 = Hybrid mask + dict",
+ " --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 types:",
+ "- [ Hash modes ] -",
"",
- "[[ Roll-your-own: Raw Hashes ]]",
+ " # | Name | Category",
+ " ------+--------------------------------------------------+--------------------------------------",
+ " 900 | MD4 | Raw Hash",
+ " 0 | MD5 | Raw Hash",
+ " 5100 | Half MD5 | Raw Hash",
+ " 100 | SHA1 | Raw Hash",
+ " 10800 | SHA-384 | Raw Hash",
+ " 1400 | SHA-256 | Raw Hash",
+ " 1700 | SHA-512 | Raw Hash",
+ " 5000 | SHA-3(Keccak) | Raw Hash",
+ " 10100 | SipHash | Raw Hash",
+ " 6000 | RipeMD160 | Raw Hash",
+ " 6100 | Whirlpool | Raw Hash",
+ " 6900 | GOST R 34.11-94 | Raw Hash",
+ " 11700 | GOST R 34.11-2012 (Streebog) 256-bit | Raw Hash",
+ " 11800 | GOST R 34.11-2012 (Streebog) 512-bit | Raw Hash",
+ " 10 | md5($pass.$salt) | Raw Hash, Salted and / or Iterated",
+ " 20 | md5($salt.$pass) | Raw Hash, Salted and / or Iterated",
+ " 30 | md5(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated",
+ " 40 | md5($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated",
+ " 3800 | md5($salt.$pass.$salt) | Raw Hash, Salted and / or Iterated",
+ " 3710 | md5($salt.md5($pass)) | Raw Hash, Salted and / or Iterated",
+ " 2600 | md5(md5($pass) | Raw Hash, Salted and / or Iterated",
+ " 4300 | md5(strtoupper(md5($pass))) | Raw Hash, Salted and / or Iterated",
+ " 4400 | md5(sha1($pass)) | Raw Hash, Salted and / or Iterated",
+ " 110 | sha1($pass.$salt) | Raw Hash, Salted and / or Iterated",
+ " 120 | sha1($salt.$pass) | Raw Hash, Salted and / or Iterated",
+ " 130 | sha1(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated",
+ " 140 | sha1($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated",
+ " 4500 | sha1(sha1($pass) | Raw Hash, Salted and / or Iterated",
+ " 4700 | sha1(md5($pass)) | Raw Hash, Salted and / or Iterated",
+ " 4900 | sha1($salt.$pass.$salt) | Raw Hash, Salted and / or Iterated",
+ " 1410 | sha256($pass.$salt) | Raw Hash, Salted and / or Iterated",
+ " 1420 | sha256($salt.$pass) | Raw Hash, Salted and / or Iterated",
+ " 1430 | sha256(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated",
+ " 1440 | sha256($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated",
+ " 1710 | sha512($pass.$salt) | Raw Hash, Salted and / or Iterated",
+ " 1720 | sha512($salt.$pass) | Raw Hash, Salted and / or Iterated",
+ " 1730 | sha512(unicode($pass).$salt) | Raw Hash, Salted and / or Iterated",
+ " 1740 | sha512($salt.unicode($pass)) | Raw Hash, Salted and / or Iterated",
+ " 50 | HMAC-MD5 (key = $pass) | Raw Hash, Authenticated",
+ " 60 | HMAC-MD5 (key = $salt) | Raw Hash, Authenticated",
+ " 150 | HMAC-SHA1 (key = $pass) | Raw Hash, Authenticated",
+ " 160 | HMAC-SHA1 (key = $salt) | Raw Hash, Authenticated",
+ " 1450 | HMAC-SHA256 (key = $pass) | Raw Hash, Authenticated",
+ " 1460 | HMAC-SHA256 (key = $salt) | Raw Hash, Authenticated",
+ " 1750 | HMAC-SHA512 (key = $pass) | Raw Hash, Authenticated",
+ " 1760 | HMAC-SHA512 (key = $salt) | Raw Hash, Authenticated",
+ " 400 | phpass | Generic KDF",
+ " 8900 | scrypt | Generic KDF",
+ " 11900 | PBKDF2-HMAC-MD5 | Generic KDF",
+ " 12000 | PBKDF2-HMAC-SHA1 | Generic KDF",
+ " 10900 | PBKDF2-HMAC-SHA256 | Generic KDF",
+ " 12100 | PBKDF2-HMAC-SHA512 | Generic KDF",
+ " 23 | Skype | Network protocols",
+ " 2500 | WPA/WPA2 | Network protocols",
+ " 4800 | iSCSI CHAP authentication, MD5(Chap) | Network protocols",
+ " 5300 | IKE-PSK MD5 | Network protocols",
+ " 5400 | IKE-PSK SHA1 | Network protocols",
+ " 5500 | NetNTLMv1 | Network protocols",
+ " 5500 | NetNTLMv1 + ESS | Network protocols",
+ " 5600 | NetNTLMv2 | Network protocols",
+ " 7300 | IPMI2 RAKP HMAC-SHA1 | Network protocols",
+ " 7500 | Kerberos 5 AS-REQ Pre-Auth etype 23 | Network protocols",
+ " 8300 | DNSSEC (NSEC3) | Network protocols",
+ " 10200 | Cram MD5 | Network protocols",
+ " 11100 | PostgreSQL CRAM (MD5) | Network protocols",
+ " 11200 | MySQL CRAM (SHA1) | Network protocols",
+ " 11400 | SIP digest authentication (MD5) | Network protocols",
+ " 13100 | Kerberos 5 TGS-REP etype 23 | Network protocols",
+ " 121 | SMF (Simple Machines Forum) | Forums, CMS, E-Commerce, Frameworks",
+ " 400 | phpBB3 | Forums, CMS, E-Commerce, Frameworks",
+ " 2611 | vBulletin < v3.8.5 | Forums, CMS, E-Commerce, Frameworks",
+ " 2711 | vBulletin > v3.8.5 | Forums, CMS, E-Commerce, Frameworks",
+ " 2811 | MyBB | Forums, CMS, E-Commerce, Frameworks",
+ " 2811 | IPB (Invison Power Board) | Forums, CMS, E-Commerce, Frameworks",
+ " 8400 | WBB3 (Woltlab Burning Board) | Forums, CMS, E-Commerce, Frameworks",
+ " 11 | Joomla < 2.5.18 | Forums, CMS, E-Commerce, Frameworks",
+ " 400 | Joomla > 2.5.18 | Forums, CMS, E-Commerce, Frameworks",
+ " 400 | Wordpress | Forums, CMS, E-Commerce, Frameworks",
+ " 2612 | PHPS | Forums, CMS, E-Commerce, Frameworks",
+ " 7900 | Drupal7 | Forums, CMS, E-Commerce, Frameworks",
+ " 21 | osCommerce | Forums, CMS, E-Commerce, Frameworks",
+ " 21 | xt:Commerce | Forums, CMS, E-Commerce, Frameworks",
+ " 11000 | PrestaShop | Forums, CMS, E-Commerce, Frameworks",
+ " 124 | Django (SHA-1) | Forums, CMS, E-Commerce, Frameworks",
+ " 10000 | Django (PBKDF2-SHA256) | Forums, CMS, E-Commerce, Frameworks",
+ " 3711 | Mediawiki B type | Forums, CMS, E-Commerce, Frameworks",
+ " 7600 | Redmine | Forums, CMS, E-Commerce, Frameworks",
+ " 12 | PostgreSQL | Database Server",
+ " 131 | MSSQL(2000) | Database Server",
+ " 132 | MSSQL(2005) | Database Server",
+ " 1731 | MSSQL(2012) | Database Server",
+ " 1731 | MSSQL(2014) | Database Server",
+ " 200 | MySQL323 | Database Server",
+ " 300 | MySQL4.1/MySQL5 | Database Server",
+ " 3100 | Oracle H: Type (Oracle 7+) | Database Server",
+ " 112 | Oracle S: Type (Oracle 11+) | Database Server",
+ " 12300 | Oracle T: Type (Oracle 12+) | Database Server",
+ " 8000 | Sybase ASE | Database Server",
+ " 141 | EPiServer 6.x < v4 | HTTP, SMTP, LDAP Server",
+ " 1441 | EPiServer 6.x > v4 | HTTP, SMTP, LDAP Server",
+ " 1600 | Apache $apr1$ | HTTP, SMTP, LDAP Server",
+ " 12600 | ColdFusion 10+ | HTTP, SMTP, LDAP Server",
+ " 1421 | hMailServer | HTTP, SMTP, LDAP Server",
+ " 101 | nsldap, SHA-1(Base64), Netscape LDAP SHA | HTTP, SMTP, LDAP Server",
+ " 111 | nsldaps, SSHA-1(Base64), Netscape LDAP SSHA | HTTP, SMTP, LDAP Server",
+ " 1711 | SSHA-512(Base64), LDAP {SSHA512} | HTTP, SMTP, LDAP Server",
+ " 11500 | CRC32 | Checksums",
+ " 3000 | LM | Operating-Systems",
+ " 1000 | NTLM | Operating-Systems",
+ " 1100 | Domain Cached Credentials (DCC), MS Cache | Operating-Systems",
+ " 2100 | Domain Cached Credentials 2 (DCC2), MS Cache 2 | Operating-Systems",
+ " 12800 | MS-AzureSync PBKDF2-HMAC-SHA256 | Operating-Systems",
+ " 1500 | descrypt, DES(Unix), Traditional DES | Operating-Systems",
+ " 12400 | BSDiCrypt, Extended DES | Operating-Systems",
+ " 500 | md5crypt $1$, MD5(Unix) | Operating-Systems",
+ " 3200 | bcrypt $2*$, Blowfish(Unix) | Operating-Systems",
+ " 7400 | sha256crypt $5$, SHA256(Unix) | Operating-Systems",
+ " 1800 | sha512crypt $6$, SHA512(Unix) | Operating-Systems",
+ " 122 | OSX v10.4, OSX v10.5, OSX v10.6 | Operating-Systems",
+ " 1722 | OSX v10.7 | Operating-Systems",
+ " 7100 | OSX v10.8, OSX v10.9, OSX v10.10 | Operating-Systems",
+ " 6300 | AIX {smd5} | Operating-Systems",
+ " 6700 | AIX {ssha1} | Operating-Systems",
+ " 6400 | AIX {ssha256} | Operating-Systems",
+ " 6500 | AIX {ssha512} | Operating-Systems",
+ " 2400 | Cisco-PIX | Operating-Systems",
+ " 2410 | Cisco-ASA | Operating-Systems",
+ " 500 | Cisco-IOS $1$ | Operating-Systems",
+ " 5700 | Cisco-IOS $4$ | Operating-Systems",
+ " 9200 | Cisco-IOS $8$ | Operating-Systems",
+ " 9300 | Cisco-IOS $9$ | Operating-Systems",
+ " 22 | Juniper Netscreen/SSG (ScreenOS) | Operating-Systems",
+ " 501 | Juniper IVE | Operating-Systems",
+ " 5800 | Android PIN | Operating-Systems",
+ " 8100 | Citrix Netscaler | Operating-Systems",
+ " 8500 | RACF | Operating-Systems",
+ " 7200 | GRUB 2 | Operating-Systems",
+ " 9900 | Radmin2 | Operating-Systems",
+ " 125 | ArubaOS | Operating-Systems",
+ " 7700 | SAP CODVN B (BCODE) | Enterprise Application Software (EAS)",
+ " 7800 | SAP CODVN F/G (PASSCODE) | Enterprise Application Software (EAS)",
+ " 10300 | SAP CODVN H (PWDSALTEDHASH) iSSHA-1 | Enterprise Application Software (EAS)",
+ " 8600 | Lotus Notes/Domino 5 | Enterprise Application Software (EAS)",
+ " 8700 | Lotus Notes/Domino 6 | Enterprise Application Software (EAS)",
+ " 9100 | Lotus Notes/Domino 8 | Enterprise Application Software (EAS)",
+ " 133 | PeopleSoft | Enterprise Application Software (EAS)",
+ " 13500 | PeopleSoft Token | Enterprise Application Software (EAS)",
+ " 11600 | 7-Zip | Archives",
+ " 12500 | RAR3-hp | Archives",
+ " 13000 | RAR5 | Archives",
+ " 13200 | AxCrypt | Archives",
+ " 13300 | AxCrypt in memory SHA1 | Archives",
+ " 13600 | WinZip | Archives",
+ " 62XY | TrueCrypt | Full-Disk encryptions (FDE)",
+ " X | 1 = PBKDF2-HMAC-RipeMD160 | Full-Disk encryptions (FDE)",
+ " X | 2 = PBKDF2-HMAC-SHA512 | Full-Disk encryptions (FDE)",
+ " X | 3 = PBKDF2-HMAC-Whirlpool | Full-Disk encryptions (FDE)",
+ " X | 4 = PBKDF2-HMAC-RipeMD160 + boot-mode | Full-Disk encryptions (FDE)",
+ " Y | 1 = XTS 512 bit pure AES | Full-Disk encryptions (FDE)",
+ " Y | 1 = XTS 512 bit pure Serpent | Full-Disk encryptions (FDE)",
+ " Y | 1 = XTS 512 bit pure Twofish | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit pure AES | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit pure Serpent | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit pure Twofish | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit cascaded AES-Twofish | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit cascaded Serpent-AES | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit cascaded Twofish-Serpent | Full-Disk encryptions (FDE)",
+ " Y | 3 = XTS 1536 bit all | Full-Disk encryptions (FDE)",
+ " 8800 | Android FDE < v4.3 | Full-Disk encryptions (FDE)",
+ " 12900 | Android FDE (Samsung DEK) | Full-Disk encryptions (FDE)",
+ " 12200 | eCryptfs | Full-Disk encryptions (FDE)",
+ " 137XY | VeraCrypt | Full-Disk encryptions (FDE)",
+ " X | 1 = PBKDF2-HMAC-RipeMD160 | Full-Disk encryptions (FDE)",
+ " X | 2 = PBKDF2-HMAC-SHA512 | Full-Disk encryptions (FDE)",
+ " X | 3 = PBKDF2-HMAC-Whirlpool | Full-Disk encryptions (FDE)",
+ " X | 4 = PBKDF2-HMAC-RipeMD160 + boot-mode | Full-Disk encryptions (FDE)",
+ " X | 5 = PBKDF2-HMAC-SHA256 | Full-Disk encryptions (FDE)",
+ " X | 6 = PBKDF2-HMAC-SHA256 + boot-mode | Full-Disk encryptions (FDE)",
+ " Y | 1 = XTS 512 bit pure AES | Full-Disk encryptions (FDE)",
+ " Y | 1 = XTS 512 bit pure Serpent | Full-Disk encryptions (FDE)",
+ " Y | 1 = XTS 512 bit pure Twofish | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit pure AES | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit pure Serpent | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit pure Twofish | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit cascaded AES-Twofish | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit cascaded Serpent-AES | Full-Disk encryptions (FDE)",
+ " Y | 2 = XTS 1024 bit cascaded Twofish-Serpent | Full-Disk encryptions (FDE)",
+ " Y | 3 = XTS 1536 bit all | Full-Disk encryptions (FDE)",
+ " 9700 | MS Office <= 2003 $0|$1, MD5 + RC4 | Documents",
+ " 9710 | MS Office <= 2003 $0|$1, MD5 + RC4, collider #1 | Documents",
+ " 9720 | MS Office <= 2003 $0|$1, MD5 + RC4, collider #2 | Documents",
+ " 9800 | MS Office <= 2003 $3|$4, SHA1 + RC4 | Documents",
+ " 9810 | MS Office <= 2003 $3|$4, SHA1 + RC4, collider #1 | Documents",
+ " 9820 | MS Office <= 2003 $3|$4, SHA1 + RC4, collider #2 | Documents",
+ " 9400 | MS Office 2007 | Documents",
+ " 9500 | MS Office 2010 | Documents",
+ " 9600 | MS Office 2013 | Documents",
+ " 10400 | PDF 1.1 - 1.3 (Acrobat 2 - 4) | Documents",
+ " 10410 | PDF 1.1 - 1.3 (Acrobat 2 - 4), collider #1 | Documents",
+ " 10420 | PDF 1.1 - 1.3 (Acrobat 2 - 4), collider #2 | Documents",
+ " 10500 | PDF 1.4 - 1.6 (Acrobat 5 - 8) | Documents",
+ " 10600 | PDF 1.7 Level 3 (Acrobat 9) | Documents",
+ " 10700 | PDF 1.7 Level 8 (Acrobat 10 - 11) | Documents",
+ " 9000 | Password Safe v2 | Password Managers",
+ " 5200 | Password Safe v3 | Password Managers",
+ " 6800 | Lastpass + Lastpass sniffed | Password Managers",
+ " 6600 | 1Password, agilekeychain | Password Managers",
+ " 8200 | 1Password, cloudkeychain | Password Managers",
+ " 11300 | Bitcoin/Litecoin wallet.dat | Password Managers",
+ " 12700 | Blockchain, My Wallet | Password Managers",
+ " 13400 | Keepass 1 (AES/Twofish) and Keepass 2 (AES) | Password Managers",
"",
- " 900 = MD4",
- " 0 = MD5",
- " 5100 = Half MD5",
- " 100 = SHA1",
- " 10800 = SHA-384",
- " 1400 = SHA-256",
- " 1700 = SHA-512",
- " 5000 = SHA-3(Keccak)",
- " 10100 = SipHash",
- " 6000 = RipeMD160",
- " 6100 = Whirlpool",
- " 6900 = GOST R 34.11-94",
- " 11700 = GOST R 34.11-2012 (Streebog) 256-bit",
- " 11800 = GOST R 34.11-2012 (Streebog) 512-bit",
+ "- [ Outfile Formats ] -",
"",
- "[[ Roll-your-own: Iterated and / or Salted Hashes ]]",
+ " # | Format",
+ " ---+--------",
+ " 1 | hash[:salt]",
+ " 2 | plain",
+ " 3 | hash[:salt]:plain",
+ " 4 | hex_plain",
+ " 5 | hash[:salt]:hex_plain",
+ " 6 | plain:hex_plain",
+ " 7 | hash[:salt]:plain:hex_plain",
+ " 8 | crackpos",
+ " 9 | hash[:salt]:crack_pos",
+ " 10 | plain:crack_pos",
+ " 11 | hash[:salt]:plain:crack_pos",
+ " 12 | hex_plain:crack_pos",
+ " 13 | hash[:salt]:hex_plain:crack_pos",
+ " 14 | plain:hex_plain:crack_pos",
+ " 15 | hash[:salt]:plain:hex_plain:crack_pos",
"",
- " 10 = md5($pass.$salt)",
- " 20 = md5($salt.$pass)",
- " 30 = md5(unicode($pass).$salt)",
- " 40 = md5($salt.unicode($pass))",
- " 3800 = md5($salt.$pass.$salt)",
- " 3710 = md5($salt.md5($pass))",
- " 2600 = md5(md5($pass)",
- " 4300 = md5(strtoupper(md5($pass)))",
- " 4400 = md5(sha1($pass))",
- " 110 = sha1($pass.$salt)",
- " 120 = sha1($salt.$pass)",
- " 130 = sha1(unicode($pass).$salt)",
- " 140 = sha1($salt.unicode($pass))",
- " 4500 = sha1(sha1($pass)",
- " 4700 = sha1(md5($pass))",
- " 4900 = sha1($salt.$pass.$salt)",
- " 1410 = sha256($pass.$salt)",
- " 1420 = sha256($salt.$pass)",
- " 1430 = sha256(unicode($pass).$salt)",
- " 1440 = sha256($salt.unicode($pass))",
- " 1710 = sha512($pass.$salt)",
- " 1720 = sha512($salt.$pass)",
- " 1730 = sha512(unicode($pass).$salt)",
- " 1740 = sha512($salt.unicode($pass))",
+ "- [ Rule Debugging Modes ] -",
"",
- "[[ Roll-your-own: Authenticated Hashes ]]",
+ " # | Format",
+ " ---+--------",
+ " 1 | Finding-Rule",
+ " 2 | Original-Word",
+ " 3 | Original-Word:Finding-Rule",
+ " 4 | Original-Word:Finding-Rule:Processed-Word",
"",
- " 50 = HMAC-MD5 (key = $pass)",
- " 60 = HMAC-MD5 (key = $salt)",
- " 150 = HMAC-SHA1 (key = $pass)",
- " 160 = HMAC-SHA1 (key = $salt)",
- " 1450 = HMAC-SHA256 (key = $pass)",
- " 1460 = HMAC-SHA256 (key = $salt)",
- " 1750 = HMAC-SHA512 (key = $pass)",
- " 1760 = HMAC-SHA512 (key = $salt)",
+ "- [ Attack Modes ] -",
"",
- "[[ Generic KDF ]]",
+ " # | Mode",
+ " ---+------",
+ " 0 | Straight",
+ " 1 | Combination",
+ " 3 | Brute-force",
+ " 6 | Hybrid Wordlist + Mask",
+ " 7 | Hybrid Mask + Wordlist",
"",
- " 400 = phpass",
- " 8900 = scrypt",
- " 11900 = PBKDF2-HMAC-MD5",
- " 12000 = PBKDF2-HMAC-SHA1",
- " 10900 = PBKDF2-HMAC-SHA256",
- " 12100 = PBKDF2-HMAC-SHA512",
+ "- [ Built-in Charsets ] -",
"",
- "[[ Network protocols, Challenge-Response ]]",
+ " ? | Charset",
+ " ---+---------",
+ " l | abcdefghijklmnopqrstuvwxyz",
+ " u | ABCDEFGHIJKLMNOPQRSTUVWXYZ",
+ " d | 0123456789",
+ " s | !\"#$%%&'()*+,-./:;<=>?@[\\]^_`{|}~",
+ " a | ?l?u?d?s",
+ " b | 0x00 - 0xff",
"",
- " 23 = Skype",
- " 2500 = WPA/WPA2",
- " 4800 = iSCSI CHAP authentication, MD5(Chap)",
- " 5300 = IKE-PSK MD5",
- " 5400 = IKE-PSK SHA1",
- " 5500 = NetNTLMv1",
- " 5500 = NetNTLMv1 + ESS",
- " 5600 = NetNTLMv2",
- " 7300 = IPMI2 RAKP HMAC-SHA1",
- " 7500 = Kerberos 5 AS-REQ Pre-Auth etype 23",
- " 8300 = DNSSEC (NSEC3)",
- " 10200 = Cram MD5",
- " 11100 = PostgreSQL Challenge-Response Authentication (MD5)",
- " 11200 = MySQL Challenge-Response Authentication (SHA1)",
- " 11400 = SIP digest authentication (MD5)",
- " 13100 = Kerberos 5 TGS-REP etype 23",
+ "- [ OpenCL Device Types ] -",
"",
- "[[ Forums, CMS, E-Commerce, Frameworks, Middleware, Wiki, Management ]]",
+ " # | Device Type",
+ " ---+-------------",
+ " 1 | CPU",
+ " 2 | GPU",
+ " 3 | FPGA, DSP, Co-Processor",
"",
- " 121 = SMF (Simple Machines Forum)",
- " 400 = phpBB3",
- " 2611 = vBulletin < v3.8.5",
- " 2711 = vBulletin > v3.8.5",
- " 2811 = MyBB",
- " 2811 = IPB (Invison Power Board)",
- " 8400 = WBB3 (Woltlab Burning Board)",
- " 11 = Joomla < 2.5.18",
- " 400 = Joomla > 2.5.18",
- " 400 = Wordpress",
- " 2612 = PHPS",
- " 7900 = Drupal7",
- " 21 = osCommerce",
- " 21 = xt:Commerce",
- " 11000 = PrestaShop",
- " 124 = Django (SHA-1)",
- " 10000 = Django (PBKDF2-SHA256)",
- " 3711 = Mediawiki B type",
- " 7600 = Redmine",
+ "- [ Workload Profiles ] -",
"",
- "[[ Database Server ]]",
+ " # | Performance | Runtime | Power Consumption | Desktop Impact",
+ " ---+-------------+---------+-------------------+----------------",
+ " 1 | Low | 2 ms | Low | Minimal",
+ " 2 | Default | 12 ms | Economic | Noticeable",
+ " 3 | High | 96 ms | High | Unresponsive",
+ " 4 | Nightmare | 480 ms | Insane | Headless",
"",
- " 12 = PostgreSQL",
- " 131 = MSSQL(2000)",
- " 132 = MSSQL(2005)",
- " 1731 = MSSQL(2012)",
- " 1731 = MSSQL(2014)",
- " 200 = MySQL323",
- " 300 = MySQL4.1/MySQL5",
- " 3100 = Oracle H: Type (Oracle 7+)",
- " 112 = Oracle S: Type (Oracle 11+)",
- " 12300 = Oracle T: Type (Oracle 12+)",
- " 8000 = Sybase ASE",
+ "If you have no idea what just happened then visit the following pages:",
"",
- "[[ HTTP, SMTP, LDAP Server ]]",
- "",
- " 141 = EPiServer 6.x < v4",
- " 1441 = EPiServer 6.x > v4",
- " 1600 = Apache $apr1$",
- " 12600 = ColdFusion 10+",
- " 1421 = hMailServer",
- " 101 = nsldap, SHA-1(Base64), Netscape LDAP SHA",
- " 111 = nsldaps, SSHA-1(Base64), Netscape LDAP SSHA",
- " 1711 = SSHA-512(Base64), LDAP {SSHA512}",
- "",
- "[[ Checksums ]]",
- "",
- " 11500 = CRC32",
- "",
- "[[ Operating-Systems ]]",
- "",
- " 3000 = LM",
- " 1000 = NTLM",
- " 1100 = Domain Cached Credentials (DCC), MS Cache",
- " 2100 = Domain Cached Credentials 2 (DCC2), MS Cache 2",
- " 12800 = MS-AzureSync PBKDF2-HMAC-SHA256",
- " 1500 = descrypt, DES(Unix), Traditional DES",
- " 12400 = BSDiCrypt, Extended DES",
- " 500 = md5crypt $1$, MD5(Unix)",
- " 3200 = bcrypt $2*$, Blowfish(Unix)",
- " 7400 = sha256crypt $5$, SHA256(Unix)",
- " 1800 = sha512crypt $6$, SHA512(Unix)",
- " 122 = OSX v10.4",
- " 122 = OSX v10.5",
- " 122 = OSX v10.6",
- " 1722 = OSX v10.7",
- " 7100 = OSX v10.8",
- " 7100 = OSX v10.9",
- " 7100 = OSX v10.10",
- " 6300 = AIX {smd5}",
- " 6700 = AIX {ssha1}",
- " 6400 = AIX {ssha256}",
- " 6500 = AIX {ssha512}",
- " 2400 = Cisco-PIX",
- " 2410 = Cisco-ASA",
- " 500 = Cisco-IOS $1$",
- " 5700 = Cisco-IOS $4$",
- " 9200 = Cisco-IOS $8$",
- " 9300 = Cisco-IOS $9$",
- " 22 = Juniper Netscreen/SSG (ScreenOS)",
- " 501 = Juniper IVE",
- " 5800 = Android PIN",
- " 8100 = Citrix Netscaler",
- " 8500 = RACF",
- " 7200 = GRUB 2",
- " 9900 = Radmin2",
- " 125 = ArubaOS",
- "",
- "[[ Enterprise Application Software (EAS) ]]",
- "",
- " 7700 = SAP CODVN B (BCODE)",
- " 7800 = SAP CODVN F/G (PASSCODE)",
- " 10300 = SAP CODVN H (PWDSALTEDHASH) iSSHA-1",
- " 8600 = Lotus Notes/Domino 5",
- " 8700 = Lotus Notes/Domino 6",
- " 9100 = Lotus Notes/Domino 8",
- " 133 = PeopleSoft",
- " 13500 = PeopleSoft Token",
- "",
- "[[ Archives ]]",
- "",
- " 11600 = 7-Zip",
- " 12500 = RAR3-hp",
- " 13000 = RAR5",
- " 13200 = AxCrypt",
- " 13300 = AxCrypt in memory SHA1",
- " 13600 = WinZip",
- "",
- "[[ Full-Disk encryptions (FDE) ]]",
- "",
- " 62XY = TrueCrypt 5.0+",
- " X = 1 = PBKDF2-HMAC-RipeMD160",
- " X = 2 = PBKDF2-HMAC-SHA512",
- " X = 3 = PBKDF2-HMAC-Whirlpool",
- " X = 4 = PBKDF2-HMAC-RipeMD160 + 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)",
- " 8800 = Android FDE < v4.3",
- " 12900 = Android FDE (Samsung DEK)",
- " 12200 = eCryptfs",
- "",
- "[[ Documents ]]",
- "",
- " 9700 = MS Office <= 2003 MD5 + RC4, oldoffice$0, oldoffice$1",
- " 9710 = MS Office <= 2003 MD5 + RC4, collider-mode #1",
- " 9720 = MS Office <= 2003 MD5 + RC4, collider-mode #2",
- " 9800 = MS Office <= 2003 SHA1 + RC4, oldoffice$3, oldoffice$4",
- " 9810 = MS Office <= 2003 SHA1 + RC4, collider-mode #1",
- " 9820 = MS Office <= 2003 SHA1 + RC4, collider-mode #2",
- " 9400 = MS Office 2007",
- " 9500 = MS Office 2010",
- " 9600 = MS Office 2013",
- " 10400 = PDF 1.1 - 1.3 (Acrobat 2 - 4)",
- " 10410 = PDF 1.1 - 1.3 (Acrobat 2 - 4) + collider-mode #1",
- " 10420 = PDF 1.1 - 1.3 (Acrobat 2 - 4) + collider-mode #2",
- " 10500 = PDF 1.4 - 1.6 (Acrobat 5 - 8)",
- " 10600 = PDF 1.7 Level 3 (Acrobat 9)",
- " 10700 = PDF 1.7 Level 8 (Acrobat 10 - 11)",
- "",
- "[[ Password Managers ]]",
- "",
- " 9000 = Password Safe v2",
- " 5200 = Password Safe v3",
- " 6800 = Lastpass",
- " 6600 = 1Password, agilekeychain",
- " 8200 = 1Password, cloudkeychain",
- " 11300 = Bitcoin/Litecoin wallet.dat",
- " 12700 = Blockchain, My Wallet",
- " 13400 = Keepass 1 (AES/Twofish) and Keepass 2 (AES)",
+ "* https://hashcat.net/wiki/#howtos_videos_papers_articles_etc_in_the_wild",
+ "* https://hashcat.net/wiki/#frequently_asked_questions",
"",
NULL
};
return exec_ms_sum / exec_ms_cnt;
}
-void status_display_automat ()
+void status_display_automate ()
{
FILE *out = stdout;
if (data.devices_status == STATUS_STARTING) return;
if (data.devices_status == STATUS_BYPASS) return;
- if (data.status_automat == 1)
+ if (data.status_automate == 1)
{
- status_display_automat ();
+ status_display_automate ();
return;
}
{
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];
}
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));
- if (device_param->vendor_id == VENDOR_ID_AMD)
+ if (device_param->device_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)
+ else if (device_param->device_vendor_id == VENDOR_ID_NV)
{
hm_device_val_to_str ((char *) fanspeed, HM_STR_BUF_SIZE, "%", hm_get_fanspeed_with_device_id (device_id));
}
#endif // HAVE_HWMON
}
+static void status_benchmark_automate ()
+{
+ 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_automate == 1)
+ {
+ status_benchmark_automate ();
+
+ 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;
{
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[hash_pos] == 1) continue;
- if (data.digests_shown[idx] == 1) continue;
+ hc_thread_mutex_lock (mux_display);
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);
+ hc_thread_mutex_unlock (mux_display);
+
+ check_hash (device_param, &cracked[i]);
}
+ myfree (cracked);
+
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_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);
}
}
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_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);
- hc_timer_get (timer, exec_time);
+ 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 };
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)
{
u32 kernel_accel = kernel_accel_min;
u32 kernel_loops = kernel_loops_min;
- // init some fake words
+ // in this case the user specified a fixed -u and -n on the commandline
+ // no way to tune anything
+ // but we need to run a few caching rounds
+
+ if ((kernel_loops_min == kernel_loops_max) && (kernel_accel_min == kernel_accel_max))
+ {
+ try_run (device_param, kernel_accel, kernel_loops);
+ try_run (device_param, kernel_accel, kernel_loops);
+ 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;
+ }
+
+ // from here it's clear we are allowed to autotune
+ // so let's init some fake words
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);
}
+ */
- // begin actual testing
+ #define VERIFIER_CNT 1
- double exec_ms_final = try_run (device_param, kernel_accel, kernel_loops);
+ // first find out highest kernel-loops that stays below target_ms
- if ((kernel_loops_min == kernel_loops_max) || (kernel_accel_min == kernel_accel_max))
+ if (kernel_loops_min < kernel_loops_max)
{
- // we do this in case the user specified a fixed -u and -n on the commandline
- // so we have a cached kernel for benchmark
-
- try_run (device_param, kernel_accel, kernel_loops);
- try_run (device_param, kernel_accel, kernel_loops);
- try_run (device_param, kernel_accel, kernel_loops);
- try_run (device_param, kernel_accel, kernel_loops);
- try_run (device_param, kernel_accel, kernel_loops);
- }
-
- // first find out highest kernel-loops that stays below target_ms
+ for (kernel_loops = kernel_loops_max; kernel_loops > kernel_loops_min; kernel_loops >>= 1)
+ {
+ double exec_ms = try_run (device_param, kernel_accel_min, kernel_loops);
- #define STEPS_CNT 10
+ for (int i = 0; i < VERIFIER_CNT; i++)
+ {
+ double exec_ms_v = try_run (device_param, kernel_accel_min, kernel_loops);
- for (kernel_loops = kernel_loops_max; kernel_loops > kernel_loops_min; kernel_loops >>= 1)
- {
- double exec_ms = try_run (device_param, kernel_accel_min, kernel_loops);
+ exec_ms = MIN (exec_ms, exec_ms_v);
+ }
- if (exec_ms < target_ms) break;
+ if (exec_ms < target_ms) break;
+ }
}
// now the same for kernel-accel but with the new kernel-loops from previous loop set
+ #define STEPS_CNT 10
+
if (kernel_accel_min < kernel_accel_max)
{
for (int i = 0; i < STEPS_CNT; i++)
double exec_ms = try_run (device_param, kernel_accel_try, kernel_loops);
- if (exec_ms > target_ms) break;
+ for (int i = 0; i < VERIFIER_CNT; i++)
+ {
+ double exec_ms_v = try_run (device_param, kernel_accel_try, kernel_loops);
+
+ exec_ms = MIN (exec_ms, exec_ms_v);
+ }
- exec_ms_final = exec_ms;
+ if (exec_ms > target_ms) break;
kernel_accel = kernel_accel_try;
}
}
- // there's a chance that we have a fixed kernel_loops but not a fixed kernel_accel
- // in such a case the above function would not create any change
- // we'll use the runtime to find out if we're allow to do last improvement
+ // at this point we want to know the actual runtime for the following reason:
+ // we need a reference for the balancing loop following up, and this
+ // the balancing loop can have an effect that the creates a new opportunity, for example:
+ // if the target is 95 ms and the current runtime is 48ms the above loop
+ // stopped the execution because the previous exec_ms was > 95ms
+ // due to the rebalance it's possible that the runtime reduces from 48ms to 47ms
+ // and this creates the possibility to double the workload -> 47 * 2 = 95ms, which is < 96ms
+
+ double exec_ms_pre_final = try_run (device_param, kernel_accel, kernel_loops);
+
+ for (int i = 0; i < VERIFIER_CNT; i++)
+ {
+ double exec_ms_pre_final_v = try_run (device_param, kernel_accel, kernel_loops);
+
+ exec_ms_pre_final = MIN (exec_ms_pre_final, exec_ms_pre_final_v);
+ }
+
+ u32 diff = kernel_loops - kernel_accel;
- if (exec_ms_final > 0)
+ if ((kernel_loops_min < kernel_loops_max) && (kernel_accel_min < kernel_accel_max))
{
- if ((exec_ms_final * 2) <= target_ms)
+ u32 kernel_accel_orig = kernel_accel;
+ u32 kernel_loops_orig = kernel_loops;
+
+ for (u32 f = 1; f < 1024; f++)
{
- const double exec_left = target_ms / exec_ms_final;
+ 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;
- const double accel_left = kernel_accel_max / kernel_accel;
+ 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++)
+ {
+ double exec_ms_v = try_run (device_param, kernel_accel_try, kernel_loops_try);
- const int exec_accel_min = MIN (exec_left, accel_left); // we want that to be int
+ exec_ms = MIN (exec_ms, exec_ms_v);
+ }
- if (exec_accel_min >= 2)
+ if (exec_ms < exec_ms_pre_final)
{
- kernel_accel *= exec_accel_min;
+ exec_ms_pre_final = exec_ms;
+
+ kernel_accel = kernel_accel_try;
+ kernel_loops = kernel_loops_try;
}
}
}
- // balancing the workload turns out to be very efficient
+ const double exec_left = target_ms / exec_ms_pre_final;
- const u32 kernel_power_balance = kernel_accel * kernel_loops;
+ const double accel_left = kernel_accel_max / kernel_accel;
- u32 sqrtv;
+ const double exec_accel_min = MIN (exec_left, accel_left); // we want that to be int
- for (sqrtv = 1; sqrtv < 0x100000; sqrtv++)
+ if (exec_accel_min >= 1.0)
{
- if ((sqrtv * sqrtv) >= kernel_power_balance) break;
- }
-
- const u32 kernel_accel_try = sqrtv;
- const u32 kernel_loops_try = sqrtv;
+ // this is safe to not overflow kernel_accel_max because of accel_left
- if ((kernel_accel_try <= kernel_accel_max) && (kernel_loops_try >= kernel_loops_min))
- {
- kernel_accel = kernel_accel_try;
- kernel_loops = kernel_loops_try;
+ kernel_accel = (double) kernel_accel * exec_accel_min;
}
- // reset fake words
+ // 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++;
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;
}
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
*/
* 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;
- 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 status_automate = STATUS_AUTOMATE;
+ 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;
-
- #define IDX_HELP 'h'
- #define IDX_VERSION 'V'
- #define IDX_VERSION_LOWER 'v'
- #define IDX_QUIET 0xff02
- #define IDX_SHOW 0xff03
- #define IDX_LEFT 0xff04
- #define IDX_REMOVE 0xff05
- #define IDX_REMOVE_TIMER 0xff37
- #define IDX_SKIP 's'
- #define IDX_LIMIT 'l'
- #define IDX_KEYSPACE 0xff35
- #define IDX_POTFILE_DISABLE 0xff06
- #define IDX_POTFILE_PATH 0xffe0
- #define IDX_DEBUG_MODE 0xff43
- #define IDX_DEBUG_FILE 0xff44
- #define IDX_INDUCTION_DIR 0xff46
- #define IDX_OUTFILE_CHECK_DIR 0xff47
- #define IDX_USERNAME 0xff07
- #define IDX_FORCE 0xff08
- #define IDX_RUNTIME 0xff09
- #define IDX_BENCHMARK 'b'
- #define IDX_HASH_MODE 'm'
- #define IDX_ATTACK_MODE 'a'
- #define IDX_RP_FILE 'r'
- #define IDX_RP_GEN 'g'
- #define IDX_RP_GEN_FUNC_MIN 0xff10
- #define IDX_RP_GEN_FUNC_MAX 0xff11
- #define IDX_RP_GEN_SEED 0xff34
- #define IDX_RULE_BUF_L 'j'
- #define IDX_RULE_BUF_R 'k'
- #define IDX_INCREMENT 'i'
- #define IDX_INCREMENT_MIN 0xff12
- #define IDX_INCREMENT_MAX 0xff13
- #define IDX_OUTFILE 'o'
- #define IDX_OUTFILE_FORMAT 0xff14
- #define IDX_OUTFILE_AUTOHEX_DISABLE 0xff39
- #define IDX_OUTFILE_CHECK_TIMER 0xff45
- #define IDX_RESTORE 0xff15
- #define IDX_RESTORE_DISABLE 0xff27
- #define IDX_STATUS 0xff17
- #define IDX_STATUS_TIMER 0xff18
- #define IDX_STATUS_AUTOMAT 0xff50
- #define IDX_LOOPBACK 0xff38
- #define IDX_WEAK_HASH_THRESHOLD 0xff42
- #define IDX_SESSION 0xff19
- #define IDX_HEX_CHARSET 0xff20
- #define IDX_HEX_SALT 0xff21
- #define IDX_HEX_WORDLIST 0xff40
- #define IDX_MARKOV_DISABLE 0xff22
- #define IDX_MARKOV_CLASSIC 0xff23
- #define IDX_MARKOV_THRESHOLD 't'
- #define IDX_MARKOV_HCSTAT 0xff24
- #define IDX_CPU_AFFINITY 0xff25
- #define IDX_OPENCL_DEVICES 'd'
- #define IDX_OPENCL_PLATFORMS 0xff72
- #define IDX_OPENCL_DEVICE_TYPES 0xff73
- #define IDX_OPENCL_VECTOR_WIDTH 0xff74
- #define IDX_WORKLOAD_PROFILE 'w'
- #define IDX_KERNEL_ACCEL 'n'
- #define IDX_KERNEL_LOOPS 'u'
- #define IDX_GPU_TEMP_DISABLE 0xff29
- #define IDX_GPU_TEMP_ABORT 0xff30
- #define IDX_GPU_TEMP_RETAIN 0xff31
- #define IDX_POWERTUNE_ENABLE 0xff41
- #define IDX_LOGFILE_DISABLE 0xff51
- #define IDX_TRUECRYPT_KEYFILES 0xff52
- #define IDX_SCRYPT_TMTO 0xff61
- #define IDX_SEGMENT_SIZE 'c'
- #define IDX_SEPARATOR 'p'
- #define IDX_BITMAP_MIN 0xff70
- #define IDX_BITMAP_MAX 0xff71
- #define IDX_CUSTOM_CHARSET_1 '1'
- #define IDX_CUSTOM_CHARSET_2 '2'
- #define IDX_CUSTOM_CHARSET_3 '3'
- #define IDX_CUSTOM_CHARSET_4 '4'
+ 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_VERSION_LOWER 'v'
+ #define IDX_QUIET 0xff02
+ #define IDX_SHOW 0xff03
+ #define IDX_LEFT 0xff04
+ #define IDX_REMOVE 0xff05
+ #define IDX_REMOVE_TIMER 0xff37
+ #define IDX_SKIP 's'
+ #define IDX_LIMIT 'l'
+ #define IDX_KEYSPACE 0xff35
+ #define IDX_POTFILE_DISABLE 0xff06
+ #define IDX_POTFILE_PATH 0xffe0
+ #define IDX_DEBUG_MODE 0xff43
+ #define IDX_DEBUG_FILE 0xff44
+ #define IDX_INDUCTION_DIR 0xff46
+ #define IDX_OUTFILE_CHECK_DIR 0xff47
+ #define IDX_USERNAME 0xff07
+ #define IDX_FORCE 0xff08
+ #define IDX_RUNTIME 0xff09
+ #define IDX_BENCHMARK 'b'
+ #define IDX_HASH_MODE 'm'
+ #define IDX_ATTACK_MODE 'a'
+ #define IDX_RP_FILE 'r'
+ #define IDX_RP_GEN 'g'
+ #define IDX_RP_GEN_FUNC_MIN 0xff10
+ #define IDX_RP_GEN_FUNC_MAX 0xff11
+ #define IDX_RP_GEN_SEED 0xff34
+ #define IDX_RULE_BUF_L 'j'
+ #define IDX_RULE_BUF_R 'k'
+ #define IDX_INCREMENT 'i'
+ #define IDX_INCREMENT_MIN 0xff12
+ #define IDX_INCREMENT_MAX 0xff13
+ #define IDX_OUTFILE 'o'
+ #define IDX_OUTFILE_FORMAT 0xff14
+ #define IDX_OUTFILE_AUTOHEX_DISABLE 0xff39
+ #define IDX_OUTFILE_CHECK_TIMER 0xff45
+ #define IDX_RESTORE 0xff15
+ #define IDX_RESTORE_DISABLE 0xff27
+ #define IDX_STATUS 0xff17
+ #define IDX_STATUS_TIMER 0xff18
+ #define IDX_STATUS_AUTOMATE 0xff50
+ #define IDX_LOOPBACK 0xff38
+ #define IDX_WEAK_HASH_THRESHOLD 0xff42
+ #define IDX_SESSION 0xff19
+ #define IDX_HEX_CHARSET 0xff20
+ #define IDX_HEX_SALT 0xff21
+ #define IDX_HEX_WORDLIST 0xff40
+ #define IDX_MARKOV_DISABLE 0xff22
+ #define IDX_MARKOV_CLASSIC 0xff23
+ #define IDX_MARKOV_THRESHOLD 't'
+ #define IDX_MARKOV_HCSTAT 0xff24
+ #define IDX_CPU_AFFINITY 0xff25
+ #define IDX_OPENCL_DEVICES 'd'
+ #define IDX_OPENCL_PLATFORMS 0xff72
+ #define IDX_OPENCL_DEVICE_TYPES 0xff73
+ #define IDX_OPENCL_VECTOR_WIDTH 0xff74
+ #define IDX_WORKLOAD_PROFILE 'w'
+ #define IDX_KERNEL_ACCEL 'n'
+ #define IDX_KERNEL_LOOPS 'u'
+ #define IDX_GPU_TEMP_DISABLE 0xff29
+ #define IDX_GPU_TEMP_ABORT 0xff30
+ #define IDX_GPU_TEMP_RETAIN 0xff31
+ #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'
+ #define IDX_BITMAP_MIN 0xff70
+ #define IDX_BITMAP_MAX 0xff71
+ #define IDX_CUSTOM_CHARSET_1 '1'
+ #define IDX_CUSTOM_CHARSET_2 '2'
+ #define IDX_CUSTOM_CHARSET_3 '3'
+ #define IDX_CUSTOM_CHARSET_4 '4'
char short_options[] = "hVvm:a:r:j:k:g:o:t:d:n:u:c:p:s:l:1:2:3:4:ibw:";
struct option long_options[] =
{
- {"help", no_argument, 0, IDX_HELP},
- {"version", no_argument, 0, IDX_VERSION},
- {"quiet", no_argument, 0, IDX_QUIET},
- {"show", no_argument, 0, IDX_SHOW},
- {"left", no_argument, 0, IDX_LEFT},
- {"username", no_argument, 0, IDX_USERNAME},
- {"remove", no_argument, 0, IDX_REMOVE},
- {"remove-timer", required_argument, 0, IDX_REMOVE_TIMER},
- {"skip", required_argument, 0, IDX_SKIP},
- {"limit", required_argument, 0, IDX_LIMIT},
- {"keyspace", no_argument, 0, IDX_KEYSPACE},
- {"potfile-disable", no_argument, 0, IDX_POTFILE_DISABLE},
- {"potfile-path", required_argument, 0, IDX_POTFILE_PATH},
- {"debug-mode", required_argument, 0, IDX_DEBUG_MODE},
- {"debug-file", required_argument, 0, IDX_DEBUG_FILE},
- {"induction-dir", required_argument, 0, IDX_INDUCTION_DIR},
- {"outfile-check-dir", required_argument, 0, IDX_OUTFILE_CHECK_DIR},
- {"force", no_argument, 0, IDX_FORCE},
- {"benchmark", no_argument, 0, IDX_BENCHMARK},
- {"restore", no_argument, 0, IDX_RESTORE},
- {"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},
- {"loopback", no_argument, 0, IDX_LOOPBACK},
- {"weak-hash-threshold",
- required_argument, 0, IDX_WEAK_HASH_THRESHOLD},
- {"session", required_argument, 0, IDX_SESSION},
- {"runtime", required_argument, 0, IDX_RUNTIME},
- {"generate-rules", required_argument, 0, IDX_RP_GEN},
- {"generate-rules-func-min",
- required_argument, 0, IDX_RP_GEN_FUNC_MIN},
- {"generate-rules-func-max",
- required_argument, 0, IDX_RP_GEN_FUNC_MAX},
- {"generate-rules-seed",
- required_argument, 0, IDX_RP_GEN_SEED},
- {"rule-left", required_argument, 0, IDX_RULE_BUF_L},
- {"rule-right", required_argument, 0, IDX_RULE_BUF_R},
- {"hash-type", required_argument, 0, IDX_HASH_MODE},
- {"attack-mode", required_argument, 0, IDX_ATTACK_MODE},
- {"rules-file", required_argument, 0, IDX_RP_FILE},
- {"outfile", required_argument, 0, IDX_OUTFILE},
- {"outfile-format", required_argument, 0, IDX_OUTFILE_FORMAT},
- {"outfile-autohex-disable",
- no_argument, 0, IDX_OUTFILE_AUTOHEX_DISABLE},
- {"outfile-check-timer",
- required_argument, 0, IDX_OUTFILE_CHECK_TIMER},
- {"hex-charset", no_argument, 0, IDX_HEX_CHARSET},
- {"hex-salt", no_argument, 0, IDX_HEX_SALT},
- {"hex-wordlist", no_argument, 0, IDX_HEX_WORDLIST},
- {"markov-disable", no_argument, 0, IDX_MARKOV_DISABLE},
- {"markov-classic", no_argument, 0, IDX_MARKOV_CLASSIC},
- {"markov-threshold", required_argument, 0, IDX_MARKOV_THRESHOLD},
- {"markov-hcstat", required_argument, 0, IDX_MARKOV_HCSTAT},
- {"cpu-affinity", required_argument, 0, IDX_CPU_AFFINITY},
- {"opencl-devices", required_argument, 0, IDX_OPENCL_DEVICES},
- {"opencl-platforms", required_argument, 0, IDX_OPENCL_PLATFORMS},
- {"opencl-device-types", required_argument, 0, IDX_OPENCL_DEVICE_TYPES},
- {"opencl-vector-width", required_argument, 0, IDX_OPENCL_VECTOR_WIDTH},
- {"workload-profile", required_argument, 0, IDX_WORKLOAD_PROFILE},
- {"kernel-accel", required_argument, 0, IDX_KERNEL_ACCEL},
- {"kernel-loops", required_argument, 0, IDX_KERNEL_LOOPS},
- {"gpu-temp-disable", no_argument, 0, IDX_GPU_TEMP_DISABLE},
+ {"help", no_argument, 0, IDX_HELP},
+ {"version", no_argument, 0, IDX_VERSION},
+ {"quiet", no_argument, 0, IDX_QUIET},
+ {"show", no_argument, 0, IDX_SHOW},
+ {"left", no_argument, 0, IDX_LEFT},
+ {"username", no_argument, 0, IDX_USERNAME},
+ {"remove", no_argument, 0, IDX_REMOVE},
+ {"remove-timer", required_argument, 0, IDX_REMOVE_TIMER},
+ {"skip", required_argument, 0, IDX_SKIP},
+ {"limit", required_argument, 0, IDX_LIMIT},
+ {"keyspace", no_argument, 0, IDX_KEYSPACE},
+ {"potfile-disable", no_argument, 0, IDX_POTFILE_DISABLE},
+ {"potfile-path", required_argument, 0, IDX_POTFILE_PATH},
+ {"debug-mode", required_argument, 0, IDX_DEBUG_MODE},
+ {"debug-file", required_argument, 0, IDX_DEBUG_FILE},
+ {"induction-dir", required_argument, 0, IDX_INDUCTION_DIR},
+ {"outfile-check-dir", required_argument, 0, IDX_OUTFILE_CHECK_DIR},
+ {"force", no_argument, 0, IDX_FORCE},
+ {"benchmark", no_argument, 0, IDX_BENCHMARK},
+ {"restore", no_argument, 0, IDX_RESTORE},
+ {"restore-disable", no_argument, 0, IDX_RESTORE_DISABLE},
+ {"status", no_argument, 0, IDX_STATUS},
+ {"status-timer", required_argument, 0, IDX_STATUS_TIMER},
+ {"status-automate", no_argument, 0, IDX_STATUS_AUTOMATE},
+ {"loopback", no_argument, 0, IDX_LOOPBACK},
+ {"weak-hash-threshold", required_argument, 0, IDX_WEAK_HASH_THRESHOLD},
+ {"session", required_argument, 0, IDX_SESSION},
+ {"runtime", required_argument, 0, IDX_RUNTIME},
+ {"generate-rules", required_argument, 0, IDX_RP_GEN},
+ {"generate-rules-func-min", required_argument, 0, IDX_RP_GEN_FUNC_MIN},
+ {"generate-rules-func-max", required_argument, 0, IDX_RP_GEN_FUNC_MAX},
+ {"generate-rules-seed", required_argument, 0, IDX_RP_GEN_SEED},
+ {"rule-left", required_argument, 0, IDX_RULE_BUF_L},
+ {"rule-right", required_argument, 0, IDX_RULE_BUF_R},
+ {"hash-type", required_argument, 0, IDX_HASH_MODE},
+ {"attack-mode", required_argument, 0, IDX_ATTACK_MODE},
+ {"rules-file", required_argument, 0, IDX_RP_FILE},
+ {"outfile", required_argument, 0, IDX_OUTFILE},
+ {"outfile-format", required_argument, 0, IDX_OUTFILE_FORMAT},
+ {"outfile-autohex-disable", no_argument, 0, IDX_OUTFILE_AUTOHEX_DISABLE},
+ {"outfile-check-timer", required_argument, 0, IDX_OUTFILE_CHECK_TIMER},
+ {"hex-charset", no_argument, 0, IDX_HEX_CHARSET},
+ {"hex-salt", no_argument, 0, IDX_HEX_SALT},
+ {"hex-wordlist", no_argument, 0, IDX_HEX_WORDLIST},
+ {"markov-disable", no_argument, 0, IDX_MARKOV_DISABLE},
+ {"markov-classic", no_argument, 0, IDX_MARKOV_CLASSIC},
+ {"markov-threshold", required_argument, 0, IDX_MARKOV_THRESHOLD},
+ {"markov-hcstat", required_argument, 0, IDX_MARKOV_HCSTAT},
+ {"cpu-affinity", required_argument, 0, IDX_CPU_AFFINITY},
+ {"opencl-devices", required_argument, 0, IDX_OPENCL_DEVICES},
+ {"opencl-platforms", required_argument, 0, IDX_OPENCL_PLATFORMS},
+ {"opencl-device-types", required_argument, 0, IDX_OPENCL_DEVICE_TYPES},
+ {"opencl-vector-width", required_argument, 0, IDX_OPENCL_VECTOR_WIDTH},
+ {"workload-profile", required_argument, 0, IDX_WORKLOAD_PROFILE},
+ {"kernel-accel", required_argument, 0, IDX_KERNEL_ACCEL},
+ {"kernel-loops", required_argument, 0, IDX_KERNEL_LOOPS},
+ {"gpu-temp-disable", no_argument, 0, IDX_GPU_TEMP_DISABLE},
#ifdef HAVE_HWMON
- {"gpu-temp-abort", required_argument, 0, IDX_GPU_TEMP_ABORT},
- {"gpu-temp-retain", required_argument, 0, IDX_GPU_TEMP_RETAIN},
+ {"gpu-temp-abort", required_argument, 0, IDX_GPU_TEMP_ABORT},
+ {"gpu-temp-retain", required_argument, 0, IDX_GPU_TEMP_RETAIN},
#ifdef HAVE_ADL
- {"powertune-enable", no_argument, 0, IDX_POWERTUNE_ENABLE},
+ {"powertune-enable", no_argument, 0, IDX_POWERTUNE_ENABLE},
#endif
#endif // HAVE_HWMON
- {"logfile-disable", no_argument, 0, IDX_LOGFILE_DISABLE},
- {"truecrypt-keyfiles", required_argument, 0, IDX_TRUECRYPT_KEYFILES},
- {"segment-size", required_argument, 0, IDX_SEGMENT_SIZE},
- {"scrypt-tmto", required_argument, 0, IDX_SCRYPT_TMTO},
- // deprecated
- {"seperator", required_argument, 0, IDX_SEPARATOR},
- {"separator", required_argument, 0, IDX_SEPARATOR},
- {"bitmap-min", required_argument, 0, IDX_BITMAP_MIN},
- {"bitmap-max", required_argument, 0, IDX_BITMAP_MAX},
- {"increment", no_argument, 0, IDX_INCREMENT},
- {"increment-min", required_argument, 0, IDX_INCREMENT_MIN},
- {"increment-max", required_argument, 0, IDX_INCREMENT_MAX},
- {"custom-charset1", required_argument, 0, IDX_CUSTOM_CHARSET_1},
- {"custom-charset2", required_argument, 0, IDX_CUSTOM_CHARSET_2},
- {"custom-charset3", required_argument, 0, IDX_CUSTOM_CHARSET_3},
- {"custom-charset4", required_argument, 0, IDX_CUSTOM_CHARSET_4},
-
+ {"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},
+ {"seperator", required_argument, 0, IDX_SEPARATOR},
+ {"separator", required_argument, 0, IDX_SEPARATOR},
+ {"bitmap-min", required_argument, 0, IDX_BITMAP_MIN},
+ {"bitmap-max", required_argument, 0, IDX_BITMAP_MAX},
+ {"increment", no_argument, 0, IDX_INCREMENT},
+ {"increment-min", required_argument, 0, IDX_INCREMENT_MIN},
+ {"increment-max", required_argument, 0, IDX_INCREMENT_MAX},
+ {"custom-charset1", required_argument, 0, IDX_CUSTOM_CHARSET_1},
+ {"custom-charset2", required_argument, 0, IDX_CUSTOM_CHARSET_2},
+ {"custom-charset3", required_argument, 0, IDX_CUSTOM_CHARSET_3},
+ {"custom-charset4", required_argument, 0, IDX_CUSTOM_CHARSET_4},
{0, 0, 0, 0}
};
{
switch (c)
{
- //case IDX_HELP: usage = 1; break;
- //case IDX_VERSION: version = 1; break;
- //case IDX_RESTORE: restore = 1; break;
- case IDX_QUIET: quiet = 1; break;
- //case IDX_SHOW: show = 1; break;
- case IDX_SHOW: break;
- //case IDX_LEFT: left = 1; break;
- case IDX_LEFT: break;
- case IDX_USERNAME: username = 1; break;
- case IDX_REMOVE: remove = 1; break;
- case IDX_REMOVE_TIMER: remove_timer = atoi (optarg);
- remove_timer_chgd = 1; break;
- case IDX_POTFILE_DISABLE: potfile_disable = 1; break;
- case IDX_POTFILE_PATH: potfile_path = optarg; break;
- case IDX_DEBUG_MODE: debug_mode = atoi (optarg); break;
- case IDX_DEBUG_FILE: debug_file = optarg; break;
- case IDX_INDUCTION_DIR: induction_dir = optarg; break;
- case IDX_OUTFILE_CHECK_DIR: outfile_check_dir = optarg; break;
- case IDX_FORCE: force = 1; break;
- case IDX_SKIP: skip = atoll (optarg); break;
- case IDX_LIMIT: limit = atoll (optarg); break;
- case IDX_KEYSPACE: keyspace = 1; break;
- case IDX_BENCHMARK: benchmark = 1; break;
- case IDX_RESTORE: break;
- 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_LOOPBACK: loopback = 1; break;
- case IDX_WEAK_HASH_THRESHOLD:
- weak_hash_threshold = atoi (optarg); break;
- //case IDX_SESSION: session = optarg; break;
- case IDX_SESSION: break;
- case IDX_HASH_MODE: hash_mode = atoi (optarg);
- hash_mode_chgd = 1; break;
- case IDX_RUNTIME: runtime = atoi (optarg);
- runtime_chgd = 1; break;
- case IDX_ATTACK_MODE: attack_mode = atoi (optarg);
- attack_mode_chgd = 1; break;
- case IDX_RP_FILE: rp_files[rp_files_cnt++] = optarg; break;
- case IDX_RP_GEN: rp_gen = atoi (optarg); break;
- case IDX_RP_GEN_FUNC_MIN: rp_gen_func_min = atoi (optarg); break;
- case IDX_RP_GEN_FUNC_MAX: rp_gen_func_max = atoi (optarg); break;
- case IDX_RP_GEN_SEED: rp_gen_seed = atoi (optarg);
- rp_gen_seed_chgd = 1; break;
- case IDX_RULE_BUF_L: rule_buf_l = optarg; break;
- case IDX_RULE_BUF_R: rule_buf_r = optarg; break;
- case IDX_MARKOV_DISABLE: markov_disable = 1; break;
- case IDX_MARKOV_CLASSIC: markov_classic = 1; break;
- case IDX_MARKOV_THRESHOLD: markov_threshold = atoi (optarg); break;
- case IDX_MARKOV_HCSTAT: markov_hcstat = optarg; break;
- case IDX_OUTFILE: outfile = optarg; break;
- case IDX_OUTFILE_FORMAT: outfile_format = atoi (optarg);
- outfile_format_chgd = 1; break;
- case IDX_OUTFILE_AUTOHEX_DISABLE:
- outfile_autohex = 0; break;
- case IDX_OUTFILE_CHECK_TIMER:
- outfile_check_timer = atoi (optarg); break;
- case IDX_HEX_CHARSET: hex_charset = 1; break;
- case IDX_HEX_SALT: hex_salt = 1; break;
- case IDX_HEX_WORDLIST: hex_wordlist = 1; break;
- case IDX_CPU_AFFINITY: cpu_affinity = optarg; break;
- case IDX_OPENCL_DEVICES: opencl_devices = optarg; break;
- case IDX_OPENCL_PLATFORMS: opencl_platforms = optarg; break;
- case IDX_OPENCL_DEVICE_TYPES:
- opencl_device_types = optarg; break;
- case IDX_OPENCL_VECTOR_WIDTH:
- opencl_vector_width = atoi (optarg);
- opencl_vector_width_chgd = 1; break;
- case IDX_WORKLOAD_PROFILE: workload_profile = atoi (optarg);
- workload_profile_chgd = 1; break;
- case IDX_KERNEL_ACCEL: kernel_accel = atoi (optarg);
- kernel_accel_chgd = 1; break;
- case IDX_KERNEL_LOOPS: kernel_loops = atoi (optarg);
- kernel_loops_chgd = 1; break;
- case IDX_GPU_TEMP_DISABLE: gpu_temp_disable = 1; break;
+ //case IDX_HELP: usage = 1; break;
+ //case IDX_VERSION: version = 1; break;
+ //case IDX_RESTORE: restore = 1; break;
+ case IDX_QUIET: quiet = 1; break;
+ //case IDX_SHOW: show = 1; break;
+ case IDX_SHOW: break;
+ //case IDX_LEFT: left = 1; break;
+ case IDX_LEFT: break;
+ case IDX_USERNAME: username = 1; break;
+ case IDX_REMOVE: remove = 1; break;
+ case IDX_REMOVE_TIMER: remove_timer = atoi (optarg);
+ remove_timer_chgd = 1; break;
+ case IDX_POTFILE_DISABLE: potfile_disable = 1; break;
+ case IDX_POTFILE_PATH: potfile_path = optarg; break;
+ case IDX_DEBUG_MODE: debug_mode = atoi (optarg); break;
+ case IDX_DEBUG_FILE: debug_file = optarg; break;
+ case IDX_INDUCTION_DIR: induction_dir = optarg; break;
+ case IDX_OUTFILE_CHECK_DIR: outfile_check_dir = optarg; break;
+ case IDX_FORCE: force = 1; break;
+ case IDX_SKIP: skip = atoll (optarg); break;
+ case IDX_LIMIT: limit = atoll (optarg); break;
+ case IDX_KEYSPACE: keyspace = 1; break;
+ case IDX_BENCHMARK: benchmark = 1; break;
+ case IDX_RESTORE: break;
+ 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_AUTOMATE: status_automate = 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;
+ case IDX_SESSION: break;
+ case IDX_HASH_MODE: hash_mode = atoi (optarg);
+ hash_mode_chgd = 1; break;
+ case IDX_RUNTIME: runtime = atoi (optarg);
+ runtime_chgd = 1; break;
+ case IDX_ATTACK_MODE: attack_mode = atoi (optarg);
+ attack_mode_chgd = 1; break;
+ case IDX_RP_FILE: rp_files[rp_files_cnt++] = optarg; break;
+ case IDX_RP_GEN: rp_gen = atoi (optarg); break;
+ case IDX_RP_GEN_FUNC_MIN: rp_gen_func_min = atoi (optarg); break;
+ case IDX_RP_GEN_FUNC_MAX: rp_gen_func_max = atoi (optarg); break;
+ case IDX_RP_GEN_SEED: rp_gen_seed = atoi (optarg);
+ rp_gen_seed_chgd = 1; break;
+ case IDX_RULE_BUF_L: rule_buf_l = optarg; break;
+ case IDX_RULE_BUF_R: rule_buf_r = optarg; break;
+ case IDX_MARKOV_DISABLE: markov_disable = 1; break;
+ case IDX_MARKOV_CLASSIC: markov_classic = 1; break;
+ case IDX_MARKOV_THRESHOLD: markov_threshold = atoi (optarg); break;
+ case IDX_MARKOV_HCSTAT: markov_hcstat = optarg; break;
+ case IDX_OUTFILE: outfile = optarg; break;
+ case IDX_OUTFILE_FORMAT: outfile_format = atoi (optarg);
+ outfile_format_chgd = 1; break;
+ case IDX_OUTFILE_AUTOHEX_DISABLE: outfile_autohex = 0; break;
+ case IDX_OUTFILE_CHECK_TIMER: outfile_check_timer = atoi (optarg); break;
+ case IDX_HEX_CHARSET: hex_charset = 1; break;
+ case IDX_HEX_SALT: hex_salt = 1; break;
+ case IDX_HEX_WORDLIST: hex_wordlist = 1; break;
+ case IDX_CPU_AFFINITY: cpu_affinity = optarg; break;
+ case IDX_OPENCL_DEVICES: opencl_devices = optarg; break;
+ case IDX_OPENCL_PLATFORMS: opencl_platforms = optarg; break;
+ case IDX_OPENCL_DEVICE_TYPES: opencl_device_types = optarg; break;
+ case IDX_OPENCL_VECTOR_WIDTH: opencl_vector_width = atoi (optarg);
+ opencl_vector_width_chgd = 1; break;
+ case IDX_WORKLOAD_PROFILE: workload_profile = atoi (optarg);
+ workload_profile_chgd = 1; break;
+ case IDX_KERNEL_ACCEL: kernel_accel = atoi (optarg);
+ kernel_accel_chgd = 1; break;
+ case IDX_KERNEL_LOOPS: kernel_loops = atoi (optarg);
+ kernel_loops_chgd = 1; break;
+ case IDX_GPU_TEMP_DISABLE: gpu_temp_disable = 1; break;
#ifdef HAVE_HWMON
- case IDX_GPU_TEMP_ABORT: gpu_temp_abort = atoi (optarg);
- #ifdef HAVE_ADL
- gpu_temp_abort_chgd = 1;
- #endif
- break;
- case IDX_GPU_TEMP_RETAIN: gpu_temp_retain = atoi (optarg);
- #ifdef HAVE_ADL
- gpu_temp_retain_chgd = 1;
- #endif
- break;
+ case IDX_GPU_TEMP_ABORT: gpu_temp_abort = atoi (optarg);
+ #ifdef HAVE_ADL
+ gpu_temp_abort_chgd = 1;
+ #endif
+ break;
+ case IDX_GPU_TEMP_RETAIN: gpu_temp_retain = atoi (optarg);
#ifdef HAVE_ADL
- case IDX_POWERTUNE_ENABLE: powertune_enable = 1; break;
+ gpu_temp_retain_chgd = 1;
+ #endif
+ break;
+ #ifdef HAVE_ADL
+ case IDX_POWERTUNE_ENABLE: powertune_enable = 1; break;
#endif
#endif // HAVE_HWMON
- case IDX_LOGFILE_DISABLE: logfile_disable = 1; break;
- case IDX_TRUECRYPT_KEYFILES: truecrypt_keyfiles = 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;
- case IDX_BITMAP_MIN: bitmap_min = atoi (optarg); break;
- case IDX_BITMAP_MAX: bitmap_max = atoi (optarg); break;
- case IDX_INCREMENT: increment = 1; break;
- case IDX_INCREMENT_MIN: increment_min = atoi (optarg);
- increment_min_chgd = 1; break;
- case IDX_INCREMENT_MAX: increment_max = atoi (optarg);
- increment_max_chgd = 1; break;
- case IDX_CUSTOM_CHARSET_1: custom_charset_1 = optarg; break;
- case IDX_CUSTOM_CHARSET_2: custom_charset_2 = optarg; break;
- case IDX_CUSTOM_CHARSET_3: custom_charset_3 = optarg; break;
- case IDX_CUSTOM_CHARSET_4: custom_charset_4 = optarg; break;
+ 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;
+ case IDX_BITMAP_MIN: bitmap_min = atoi (optarg); break;
+ case IDX_BITMAP_MAX: bitmap_max = atoi (optarg); break;
+ case IDX_INCREMENT: increment = 1; break;
+ case IDX_INCREMENT_MIN: increment_min = atoi (optarg);
+ increment_min_chgd = 1; break;
+ case IDX_INCREMENT_MAX: increment_max = atoi (optarg);
+ increment_max_chgd = 1; break;
+ case IDX_CUSTOM_CHARSET_1: custom_charset_1 = optarg; break;
+ case IDX_CUSTOM_CHARSET_2: custom_charset_2 = optarg; break;
+ case IDX_CUSTOM_CHARSET_3: custom_charset_3 = optarg; break;
+ case IDX_CUSTOM_CHARSET_4: custom_charset_4 = optarg; break;
default:
log_error ("ERROR: Invalid argument specified");
{
if (benchmark == 1)
{
- log_info ("%s (%s) starting in benchmark-mode...", PROGNAME, VERSION_TAG);
- log_info ("");
+ if (status_automate == 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)
{
return (-1);
}
- if (hash_mode_chgd && hash_mode > 13600) // 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));
if (kernel_accel_chgd == 1)
{
+ if (force == 0)
+ {
+ log_info ("The manual use of the option -n (or --kernel-accel) is outdated");
+ log_info ("Please consider using the option -w instead");
+ log_info ("You can use --force to override this but do not post error reports if you do so");
+ log_info ("");
+
+ return (-1);
+ }
+
if (kernel_accel < 1)
{
log_error ("ERROR: Invalid kernel-accel specified");
if (kernel_loops_chgd == 1)
{
+ if (force == 0)
+ {
+ log_info ("The manual use of the option -u (or --kernel-loops) is outdated");
+ log_info ("Please consider using the option -w instead");
+ log_info ("You can use --force to override this but do not post error reports if you do so");
+ log_info ("");
+
+ return (-1);
+ }
+
if (kernel_loops < 1)
{
log_error ("ERROR: Invalid kernel-loops specified");
}
}
- 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;
- #if defined(HAVE_HWMON) && defined(HAVE_ADL)
- data.powertune_enable = powertune_enable;
+ 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_automate = status_automate;
+ 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;
+ #endif
#endif
- data.logfile_disable = logfile_disable;
+ data.logfile_disable = logfile_disable;
data.truecrypt_keyfiles = truecrypt_keyfiles;
- data.scrypt_tmto = scrypt_tmto;
- data.workload_profile = workload_profile;
+ 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 (status_automate);
logfile_top_uint (status_timer);
logfile_top_uint (usage);
logfile_top_uint (username);
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;
case 13400: hash_type = HASH_TYPE_AES;
salt_type = SALT_TYPE_EMBEDDED;
attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
- opts_type = OPTS_TYPE_PT_GENERATE_LE;
- kern_type = KERN_TYPE_KEEPASS;
- dgst_size = DGST_SIZE_4_4;
- parse_func = keepass_parse_hash;
- sort_by_digest = sort_by_digest_4_4;
+ opts_type = OPTS_TYPE_PT_GENERATE_LE;
+ kern_type = KERN_TYPE_KEEPASS;
+ dgst_size = DGST_SIZE_4_4;
+ parse_func = keepass_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;
+
+ case 13500: hash_type = HASH_TYPE_SHA1;
+ salt_type = SALT_TYPE_EMBEDDED;
+ attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
+ opts_type = OPTS_TYPE_PT_GENERATE_BE
+ | OPTS_TYPE_PT_UNICODE
+ | OPTS_TYPE_PT_ADD80;
+ kern_type = KERN_TYPE_PSTOKEN;
+ dgst_size = DGST_SIZE_4_5;
+ parse_func = pstoken_parse_hash;
+ sort_by_digest = sort_by_digest_4_5;
+ opti_type = OPTI_TYPE_ZERO_BYTE
+ | OPTI_TYPE_PRECOMPUTE_INIT
+ | OPTI_TYPE_EARLY_SKIP
+ | OPTI_TYPE_NOT_ITERATED
+ | OPTI_TYPE_PREPENDED_SALT
+ | OPTI_TYPE_RAW_HASH;
+ dgst_pos0 = 3;
+ dgst_pos1 = 4;
+ dgst_pos2 = 2;
+ dgst_pos3 = 1;
+ break;
+
+ 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;
+
+ 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;
+
+ 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;
+
+ 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;
+
+ 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;
+
+ 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;
+
+ 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;
+
+ 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_pos3 = 3;
break;
- case 13500: hash_type = HASH_TYPE_SHA1;
+ case 13762: hash_type = HASH_TYPE_SHA256;
salt_type = SALT_TYPE_EMBEDDED;
- attack_exec = ATTACK_EXEC_INSIDE_KERNEL;
- opts_type = OPTS_TYPE_PT_GENERATE_BE
- | OPTS_TYPE_PT_UNICODE
- | OPTS_TYPE_PT_ADD80;
- kern_type = KERN_TYPE_PSTOKEN;
- dgst_size = DGST_SIZE_4_5;
- parse_func = pstoken_parse_hash;
- sort_by_digest = sort_by_digest_4_5;
- opti_type = OPTI_TYPE_ZERO_BYTE
- | OPTI_TYPE_PRECOMPUTE_INIT
- | OPTI_TYPE_EARLY_SKIP
- | OPTI_TYPE_NOT_ITERATED
- | OPTI_TYPE_PREPENDED_SALT
- | OPTI_TYPE_RAW_HASH;
- dgst_pos0 = 3;
- dgst_pos1 = 4;
+ 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 = 1;
+ dgst_pos3 = 3;
break;
- case 13600: hash_type = HASH_TYPE_PBKDF2_SHA1;
+ 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;
- kern_type = KERN_TYPE_ZIP2;
- dgst_size = DGST_SIZE_4_5;
- parse_func = zip2_parse_hash;
- sort_by_digest = sort_by_digest_4_5;
+ 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_pos3 = 3;
break;
+
default: usage_mini_print (PROGNAME); return (-1);
}
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;
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;
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;
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 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;
// 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_GENERIC;
+ platform_vendor_id = VENDOR_ID_APPLE;
}
else if (strcmp (platform_vendor, CL_VENDOR_INTEL_BEIGNET) == 0)
{
- vendor_id = VENDOR_ID_GENERIC;
+ platform_vendor_id = VENDOR_ID_INTEL_BEIGNET;
}
else if (strcmp (platform_vendor, CL_VENDOR_INTEL_SDK) == 0)
{
- vendor_id = VENDOR_ID_GENERIC;
+ platform_vendor_id = VENDOR_ID_INTEL_SDK;
}
else if (strcmp (platform_vendor, CL_VENDOR_MESA) == 0)
{
- vendor_id = VENDOR_ID_GENERIC;
+ 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_GENERIC;
+ 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 (device_param->skipped == 0)
+ if (status_automate == 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 (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))
{
- log_info ("");
+ if (status_automate == 0)
+ {
+ log_info ("");
+ }
}
/**
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));
}
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 (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;
}
* 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)
{
- device_param->kernel_loops_max = data.salts_buf[0].salt_iter;
+ 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
+ {
+ innerloop_cnt = data.salts_buf[0].salt_iter;
}
- }
-
- /**
- * some algorithms need a special kernel-accel
- */
-
- if (hash_mode == 8900)
- {
- device_param->kernel_accel_min = 1;
- device_param->kernel_accel_max = 64;
- }
- if (hash_mode == 9300)
- {
- device_param->kernel_accel_min = 1;
- device_param->kernel_accel_max = 64;
+ if ((innerloop_cnt >= device_param->kernel_loops_min) &&
+ (innerloop_cnt <= device_param->kernel_loops_max))
+ {
+ device_param->kernel_loops_max = innerloop_cnt;
+ }
}
u32 kernel_accel_min = device_param->kernel_accel_min;
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
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);
+
+ 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
+
+ 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
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)
* 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;
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
*/
+ 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]);
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_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;
+
+ // 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 initial fanspeed if gpu_temp_retain is enabled
*/
if (benchmark == 1)
{
- quiet = 0;
+ if (status_automate == 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 ("");
+ }
}
/**
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);
{
status_benchmark ();
- log_info ("");
+ if (status_automate == 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);
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