Commit Graph

1629 Commits (376baa2b1005a666eb93a672ec5f8c5aa96d47e4)

Author SHA1 Message Date
jsteube 503f676cb0 Remove static keywords from OpenCL functions, they can cause old NV JiT compiler to fail
7 years ago
jsteube 9c832092df Fix some function declarations inside LUKS kernel
7 years ago
jsteube b4c3df876c Fix some formating in -m 600
7 years ago
jsteube 7c2dadab17 no volatile required here
7 years ago
Jens Steube 55eaff3a45 Merge pull request #1365 from Fist0urs/DPAPI_reworked
7 years ago
Fist0urs e3cb3e9b4c test.pl ready and 0 error. Ready for PR
7 years ago
jsteube 7a17b8159b all() function is not working as expected in scalar datatype case
7 years ago
Fist0urs a6294537fd Splitted DPAPI kernel in 2 to increase performances
7 years ago
jsteube 207ce9b853 all() function is not working as expected in scalar datatype case
7 years ago
jsteube ddbe805c00 Fix last step of make_kn
7 years ago
jsteube 68f5b12754 Get rid of swap32() in make_kn() in -m 2500
7 years ago
jsteube ab1dabebbe Fix missing include in -m 2501
7 years ago
jsteube b14f44dcf7 Fix uninitialized keymic buffer
7 years ago
jsteube beab5457e6 Backport WPA-PSK-SHA256-AES-CMAC to -m 2501
7 years ago
jsteube ca1b6492e7 Some code simplify on AES CMAC
7 years ago
jsteube 4e3a642f7f Initial WPA2-PSK-SHA256-AES-CMAC support
7 years ago
jsteube 98fc02e04b Add PTK compute for keyver 3
7 years ago
jsteube 617dbb97ba Prepare migration -m 15800 into -m 2500
7 years ago
philsmd 2dadae4e9a fixed incorrect use of the esalt_bufs for -m 600 = BLAKE2-512
7 years ago
mhasbini de7ccd88ef Fix overflow in mangle_dupechar_last function
7 years ago
jsteube b169653b8f Fix missing return value in rule_op_mangle_toggle_at()
7 years ago
jsteube 55f653f374 Get rid of volatile in TrueCrypt kernels
7 years ago
jsteube 16e33b20fc Fix out of boundary access in -m 4700
7 years ago
jsteube 51dd982b12 Bring back some volatile for AMD
7 years ago
jsteube 9125062ffc Move volatiles for AMD closer to the problem
7 years ago
jsteube 1963b12fdc According to AMD docs, GCN 3 and 4 are the same
7 years ago
jsteube ac9f1da747 Add fine-tuned AMD GCN control macros
7 years ago
Jens Steube a910aea9e0 Do not use __local memory for -m 13100 if running on a device without physical shared memory
7 years ago
Jens Steube b58aa445b4 Do not use __local memory for -m 7500 if running on a device without physical shared memory
7 years ago
Jens Steube 63f6ca5114 Do not use __local memory for whirlpool if running on a device without physical shared memory
7 years ago
jsteube 8b0e7087c7 Fixed an invalid optimization code in kernel 7700 depending on the input hash, causing the kernel to loop forever
7 years ago
jsteube 151dbc5349 Fix replace value in inc_hash_ripemd160.cl
7 years ago
jsteube f859f466ef Fix -m 8300 in -a 0 mode
7 years ago
jsteube f5e04254dc Fix -m 10800 in -a 0 mode
7 years ago
jsteube d3b9febb30 Fix some double variable declarations
7 years ago
jsteube 40b57677cd OpenCL Kernels: Reactivate Dalibors XOR optimization on MD5_H on all MD5 based algorithms
7 years ago
jsteube 6d112aeb39 OpenCL Kernels: Rewritten Keccak kernel to run fully on registers and partially reversed last round
7 years ago
jsteube a378abee66 Add missing NEW_SIMD_CODE in -m 6600
7 years ago
jsteube 1c169af0ad Make -m 14100 a pure kernel only
7 years ago
jsteube 2b9888486e Make -m 14000 a pure kernel only and add volatile for asm statement
7 years ago
jsteube 99f416435e Fix invalid use of __constant in LM kernel
7 years ago
jsteube 6db2f4cc18 Fix typo
7 years ago
jsteube 918578bee1 Improve some NVidia specific inline assembly
7 years ago
jsteube 9de1e557bb More VEGA specific inline assembly to improve SHA1 based kernels
7 years ago
jsteube a0be36d7b8 Fix compile error caused by __add3()
7 years ago
jsteube 00e38cc2c6 Add VEGA specific inline assembly to improve all MD4, MD5, SHA1 and SHA256 based kernels
7 years ago
jsteube 7bfd343ec9 Optimized rule_op_mangle_dupechar_last(), rule_op_mangle_rotate_right(), rule_op_mangle_rotate_left() and append_block1() in rule engine
7 years ago
jsteube 52a97fee75 Improve rule engine performance by improving append_0x80_xxx() performance by using precomputed values from constant memory
7 years ago
jsteube 3260000357 Fix whirlpool pure kernel in -a 0 mode
7 years ago
jsteube e3810d054b Fix some use of pw_t tmp variable
7 years ago
jsteube 5e01ff4c53 Refactor some u32x to u32 where u32x is not needed
7 years ago
jsteube 1aa76eac15 Refactor use of __constant to match up with the user selected attack mode
7 years ago
jsteube 938c281ee0 Resurrect some volatile variables in order to correctly compile pure kernels on AMD drivers
7 years ago
jsteube 48fbe81a09 Add more inline assembly for AMD ROCm
7 years ago
jsteube 6c619155c3 Workaround ROCm compiler error in aes256_ExpandKey()
7 years ago
jsteube 8c9c36ee2a Fix out-of-bound access in aesXXX_InvertKey()
7 years ago
jsteube bed7e8f466 Remove unused truncate_block_xxx_xx() functions and update kernels to use the _S function
7 years ago
jsteube 51dc1c7db3 Use truncate_block_4x4_le_S() instead of truncate_block_4x4_le() in -m 6800
7 years ago
jsteube 9b73c464d2 Fix typo in macro
7 years ago
jsteube 7b443ee7ff Optimize performance of rule_op_mangle_title_sep(), rule_op_mangle_purgechar() and rule_op_mangle_replace()
7 years ago
jsteube 0de41c2716 Some more optimizations for rule engine
7 years ago
jsteube 9f8c5a253d More rule engine performance optimizations
7 years ago
jsteube 0783289e2f Optimized a0 pure kernel for AMD
7 years ago
jsteube a5659d5619 Also switch optimized kernels rule engine to make use of kernel rules in constant memory
7 years ago
jsteube 1d04de3a8e Limit kernel-loops in straight-mode to 256, therefore allow rules to be stored in constant memory
7 years ago
jsteube 51372438fe Allow OpenCL kernel inline assembly if ROCm drivers was detected
7 years ago
jsteube 8853884f2a Fix append_four_byte() in case sm8 is 0
7 years ago
jsteube f32e113942 Add missing case in append_block() in pure kernel rule engine
7 years ago
jsteube 6907981f08 Backport current state of optimized kernel rule engine to CPU
7 years ago
jsteube 508f1562f2 Fix --stdout kernels, gid_max was still set to u32
7 years ago
jsteube 319799bbbf Switch the datatypes of the variables responsible for work-item count and work-item size from u32 to u64
7 years ago
jsteube d9c906e134 Move 0x80 to hardcoded position for sha3-256 bit in order to allow ROCm compiler to use registers only
7 years ago
jsteube 694cc0b740 Remove all calls to overwrite_at_* functions
7 years ago
jsteube e984a829ea Remove no longer needed overwrite_at_* functions
7 years ago
jsteube bf299fe043 Optimized 3DES for rocm
7 years ago
jsteube ad1ce462d1 Get rid of ceil() in OpenCL kernels
7 years ago
jsteube 53f53fe014 Reduced number of required registers in SIP based on maximum possible esalt length
7 years ago
jsteube 9ee5da40e0 Workaround rocm compiler error for -m 15300
7 years ago
jsteube 88e995ddcf Replace some SIMD related function calls
7 years ago
jsteube 5b5bdf3889 Replace some SIMD related function calls
7 years ago
jsteube 967e96728d Make all the OpenCL kernel function includes static
7 years ago
jsteube 21e9c63d46 Fix rotl64() the same was as rotr64()
7 years ago
jsteube 58012ada0c Fall back to old rotr64 optimization for AMD
7 years ago
philsmd 4a89172140
reformatting; replaced some tabs with spaces
7 years ago
jsteube ec874c1d59 Optimized the following pure kernel rule engine functions:
7 years ago
jsteube 9a57c4b20e Fix missing boundary check in pure kernel rule engine
7 years ago
jsteube dcaa91a88f Fix rule engine function call in amp_a0
7 years ago
jsteube 0b0abb5c12 Prepare pure kernel rule engine for performance optimization
7 years ago
jsteube 98b4aab9d0 Update inc_rp_optimized to inc_rp for pure kernels
7 years ago
jsteube cb1fe251de Add missing functions to pure kernel rule engine
7 years ago
jsteube 07b54c1257 Replace code to use pure kernel rule engine for slow hashes
7 years ago
jsteube 8a0d21360b Prepare for pure kernel rule engine support
7 years ago
jsteube 34d882a116 Rename inc_rp.X to inc_rp_optimized.X
7 years ago
jsteube 68a8f70edb Mix in pure kernel functions in various optimized kernels
7 years ago
jsteube 66f7590883 Mix in pure kernel functions in various optimized kernels
7 years ago
jsteube a228e296da Fix some whirlpool vector function declaration
7 years ago
jsteube 560a786ea7 Mix in pure kernel functions in various optimized kernels
7 years ago
jsteube 4b2f3011e9 Mix in pure kernel functions to PDF 1.7 Level 8 (Acrobat 10 - 11)
7 years ago
jsteube a956a84edb Fix vector datatype in -m 2810
7 years ago
jsteube e6cb69e4cb Add pure kernels for JKS Java Key Store Private Keys (SHA1)
7 years ago
jsteube adacccecdf Add pure kernels for FileZilla Server >= 0.9.55
7 years ago
jsteube 4c7f61e473 Add pure kernels for sha1(CX)
7 years ago
jsteube b4c9f46205 Prepare sha1(CX) optimized kernel for pure kernel version
7 years ago
jsteube 29e13d6b77 Add pure kernels for OpenCart
7 years ago
jsteube b6cf3144de Prepare OpenCart optimized kernel for pure kernel version
7 years ago
jsteube 4443ecd861 Add pure kernels for Windows Phone 8+ PIN/password
7 years ago
jsteube a5c0aa6041 Add pure kernels for PeopleSoft PS_TOKEN
7 years ago
jsteube 24a2fb01aa Fix missing barrier in -m 8500
7 years ago
jsteube 0a676b549f Remove global barrier when not needed to workaround Intel OpenCL runtime bug
7 years ago
jsteube b9876c100b Add pure kernels for AxCrypt in-memory SHA1
7 years ago
jsteube 5c6b3fa7ab Add pure kernels for Kerberos 5 TGS-REP etype 23
7 years ago
jsteube 51128473bc Add pure kernels for ColdFusion 10+
7 years ago
jsteube 4f72c8bee6 Add pure kernels for SIP digest authentication (MD5)
7 years ago
jsteube b1f9ed4a7c Add pure kernels for MySQL CRAM (SHA1)
7 years ago
jsteube 7548e5f85a Add pure kernels for PostgreSQL CRAM (MD5)
7 years ago
jsteube c5c12f89c1 Rewrite code around amd_bytealign to be of type BE to save a branch afterwards
7 years ago
jsteube a53d9e09de Fix some issue with offset_minus_4
7 years ago
jsteube c9cae1f663 Add pure kernels for PrestaShop
7 years ago
jsteube bc9f721dcd Mix in pure kernel functions to PDF 1.1 - 1.3 (Acrobat 2 - 4)
7 years ago
jsteube 177800d1d0 Add pure kernels for RAdmin2
7 years ago
jsteube 3a042972a4 Mix in pure kernel functions to MS Office <= 2003 $3/$4, SHA1 + RC4
7 years ago
jsteube a650b0864e Mix in pure kernel functions to MS Office <= 2003 $0/$1, MD5 + RC4
7 years ago
jsteube 37432b19bc Mix in pure kernel functions to MS Office <= 2003 $0/$1, MD5 + RC4
7 years ago
jsteube 14983a7542 Simplify RACF kernel
7 years ago
jsteube 1f42377931 Simplify Lotus Notes/Domino 5 kernel
7 years ago
jsteube c68191e47a Add pure kernels for DNSSEC (NSEC3)
7 years ago
jsteube 54eb0b158d Prepare DNSSEC (NSEC3) optimized kernel for pure kernel version
7 years ago
jsteube c2a8ae0207 Add pure kernels for WBB3 (Woltlab Burning Board)
7 years ago
jsteube 7cf3c29ef5 Mix in pure kernel functions to DNSSEC (NSEC3)
7 years ago
jsteube aafda5fa1b Add pure kernels for Citrix NetScaler
7 years ago
jsteube 5da64a1a43 Mix in pure kernel functions to SAP CODVN F/G (PASSCODE)
7 years ago
jsteube bc6b8ca1c9 Mix in pure kernel functions to SAP CODVN B (BCODE)
7 years ago
jsteube 89d52f8209 Add pure kernels for Kerberos 5 AS-REQ Pre-Auth etype 23
7 years ago
jsteube af6052d34b Revert some invalid rename of kernel files
7 years ago
jsteube 2802f1d592 Fix vector function calls and datatypes
7 years ago
jsteube 1eb249c5b4 Add pure kernels for IPMI2 RAKP HMAC-SHA1
7 years ago
jsteube 83d37ebeff Add pure kernels for FortiGate (FortiOS)
7 years ago
jsteube cbd8f81a1c Add pure kernels for RipeMD160
7 years ago
jsteube 6946329b02 Fix BF pure kernels for NetNTLMv2
7 years ago
jsteube a9fed50ce0 Add pure kernels for NetNTLMv2
7 years ago
jsteube 443fa960d3 Fix vector function calls
7 years ago
jsteube 50aeade65c Add pure kernels for NetNTLMv1 / NetNTLMv1+ESS
7 years ago
jsteube 5d137ba036 Add pure kernels for IKE-PSK SHA1
7 years ago
jsteube 942b7068be Add pure kernels for IKE-PSK MD5
7 years ago
jsteube b541e46b9b Add pure kernels for Half MD5
7 years ago
jsteube e0a565234a Optimized -m 7700 for ROCm
7 years ago
jsteube 332396a003 Fix SCRYPT on ROCm
7 years ago
jsteube a85be1d0f0 Fix some const keywords in inc_truecrypt_xts.cl
7 years ago
jsteube 02e2279d59 Optimized -m 8500 for ROCm
7 years ago
jsteube 5bcda7d05a Optimized -m 5300 and -m 5400 for ROCm
7 years ago
jsteube 772441448a Optimized -m 8000 for ROCm
7 years ago
jsteube 9562d07264 Replace bitwise swaps with rotate() versions for AMD
7 years ago
jsteube 3125a756d9 Remove some AMD _unroll restrictions no longer required with ROCm
7 years ago
jsteube 4dca908cdf Fix a typo in OpenCL/m01460_a3-optimized.cl
7 years ago
jsteube 4c71bc984e Fix const keywords in -m 8600
7 years ago
jsteube c255a967df Fix some types in rotate functions
7 years ago
jsteube c289fcb2b0 Merge branch 'master' of https://github.com/hashcat/hashcat
7 years ago
jsteube 5e34ec348e Optimize kernels for ROCm 1.6
7 years ago
jsteube 27edc07c2f Add pure kernels for iSCSI CHAP authentication, MD5(CHAP)
7 years ago
jsteube e9821a01ba Add pure kernels for sha1($salt.$pass.$salt)
7 years ago
jsteube 9515927cf7 Add pure kernels for sha1(md5())
7 years ago
jsteube 1fdb9d1d7e Add pure kernels for sha1($salt.sha1($pass))
7 years ago
jsteube 15d725e6cc Add pure kernels for sha1(sha1($pass))
7 years ago
jsteube 4e97a4db24 Add pure kernels for md5(sha1($pass))
7 years ago
jsteube 03bb234045 Preparation for WPA/WPA2 AES-CMAC: works till PMK
7 years ago
jsteube 441434840c Fix broken -m 7900 after migration to pure kernel
7 years ago
jsteube eae9329761 Workaround some AMD JiT compiler segfault on complex kernels
7 years ago
jsteube 920911b56e Migrate MD5-PIX and MD5-ASA to run as optimized kernels
7 years ago
jsteube beb6ee2061 Add OPTI_TYPE_OPTIMIZED_KERNEL
7 years ago
jsteube fbea72ebd6 Renamed default kernels to optimized kernels
7 years ago
jsteube 2616d54794 Add -L kernel for md5(strtoupper(md5($pass)))
7 years ago
jsteube 0e7d40dd3d Add -L kernel for md5($salt.md5($pass.$salt))
7 years ago
jsteube ad9ee5c5b1 Add -L kernel for md5($salt.md5($salt.$pass))
7 years ago
jsteube 80324918e9 Add -L kernel for md5(md5($pass).md5($salt))
7 years ago
jsteube 945cf9be2f md5($salt.$pass.$salt)
7 years ago
jsteube 50760b57e4 Add -L kernel for md5(.md5()), MediaWiki B type
7 years ago
jsteube d08c7689bb Add -L kernel for md5(md5(s).md5(p)), MyBB 1.2+, IPB2+
7 years ago
jsteube 90a8f58459 Add -L kernel for md5(md5()), vBulletin, PHPS
7 years ago
jsteube e863a12624 Some fixes for inc_hash_sha224.cl and inc_hash_sha384.cl
7 years ago
jsteube fe38379d0d Add -L kernel for SHA224
7 years ago
jsteube dbf74b68b2 Add -L kernel for Domain Cached Credentials (DCC), MS Cache
7 years ago
jsteube cc4ff214d6 Add -L kernel for NTLM
7 years ago
jsteube b2f3bfb06c Add -L kernel for MD4
7 years ago
Jens Steube 3d9b071e1e Improve CPU cracking speed by replacing vector comparison functions with the more advanced ones available on CPU
7 years ago
jsteube 2dd1833998 Move from ld.global.v4.u32 to ld.const.v4.u32 in _a3 kernels
7 years ago
jsteube cd313c9c28 Add -L kernel for MySQL4.1/MySQL5
7 years ago
jsteube 757f3a39c2 Accidentially pushed experimental -m 2500 kernel
7 years ago
jsteube 8434e451ef Add -L support for all SHA512 based generic hashes
7 years ago
jsteube 7205f450dd Backport more HMAC functions in inc_hash_xxx.cl from global to private
7 years ago
jsteube 0334cdb277 Fix some datatypes
7 years ago
jsteube 5c75eb84b0 Add -L support for all SHA256 based generic hashes
7 years ago
jsteube e2371540e0 Add missing kernel in -m m00150_a1-pure.cl and m00160_a1-pure.cl
7 years ago
jsteube c4098e2230 Fix invalid use of a non-vector function from within a vector function
7 years ago
jsteube e4683aebb8 Add -L support for all sha1 based generic hashes
7 years ago
jsteube 696afc2a1b Fix selector in switch_buffer_by_offset_1x64_be_S()
7 years ago
jsteube 4e0972ce3a Add xxx_update_vector_swap(), xxx_update_vector_utf16le_swap() and xxx_update_vector_utf16beN() for later use
7 years ago
jsteube 9c6c21490f Add *_hmac_init_swap for later use
7 years ago
jsteube f03156b05e Add switch_buffer_by_offset_1x64_be_S() and code generators for later use
7 years ago
jsteube 5707fed499 Add example -L kernel for algorithms using HMAC where the password is the data
7 years ago
jsteube 994e7efc91 Add example -L kernel for algorithms using HMAC where the password is the key
7 years ago