diff --git a/OpenCL/m02400_a0-optimized.cl b/OpenCL/m02400_a0-optimized.cl index bd3a613ea..be29db7b3 100644 --- a/OpenCL/m02400_a0-optimized.cl +++ b/OpenCL/m02400_a0-optimized.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -57,18 +57,39 @@ __kernel void m02400_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); - w1[0] = 0x80; - w1[1] = 0; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 16 * 8; - w3[3] = 0; + if (out_len <= 16) + { + w1[0] = 0x80; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 16 * 8; + w3[3] = 0; + } + else if (out_len <= 32) + { + w2[0] = 0x80; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 32 * 8; + w3[3] = 0; + } + else if (out_len <= 48) + { + w3[0] = 0x80; + w3[1] = 0; + w3[2] = 48 * 8; + w3[3] = 0; + } u32x a = MD5M_A; u32x b = MD5M_B; @@ -217,18 +238,39 @@ __kernel void m02400_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); - w1[0] = 0x80; - w1[1] = 0; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 16 * 8; - w3[3] = 0; + if (out_len <= 16) + { + w1[0] = 0x80; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 16 * 8; + w3[3] = 0; + } + else if (out_len <= 32) + { + w2[0] = 0x80; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 32 * 8; + w3[3] = 0; + } + else if (out_len <= 48) + { + w3[0] = 0x80; + w3[1] = 0; + w3[2] = 48 * 8; + w3[3] = 0; + } u32x a = MD5M_A; u32x b = MD5M_B; diff --git a/OpenCL/m02400_a1-optimized.cl b/OpenCL/m02400_a1-optimized.cl index 40481dea2..30b54c509 100644 --- a/OpenCL/m02400_a1-optimized.cl +++ b/OpenCL/m02400_a1-optimized.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -119,18 +119,39 @@ __kernel void m02400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule * md5 */ - w1[0] = 0x80; - w1[1] = 0; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 16 * 8; - w3[3] = 0; + if (pw_len <= 16) + { + w1[0] = 0x80; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 16 * 8; + w3[3] = 0; + } + else if (pw_len <= 32) + { + w2[0] = 0x80; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 32 * 8; + w3[3] = 0; + } + else if (pw_len <= 48) + { + w3[0] = 0x80; + w3[1] = 0; + w3[2] = 48 * 8; + w3[3] = 0; + } u32x a = MD5M_A; u32x b = MD5M_B; @@ -343,18 +364,39 @@ __kernel void m02400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule * md5 */ - w1[0] = 0x80; - w1[1] = 0; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 16 * 8; - w3[3] = 0; + if (pw_len <= 16) + { + w1[0] = 0x80; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 16 * 8; + w3[3] = 0; + } + else if (pw_len <= 32) + { + w2[0] = 0x80; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 32 * 8; + w3[3] = 0; + } + else if (pw_len <= 48) + { + w3[0] = 0x80; + w3[1] = 0; + w3[2] = 48 * 8; + w3[3] = 0; + } u32x a = MD5M_A; u32x b = MD5M_B; diff --git a/OpenCL/m02400_a3-optimized.cl b/OpenCL/m02400_a3-optimized.cl index bd29d259c..0a411e3d3 100644 --- a/OpenCL/m02400_a3-optimized.cl +++ b/OpenCL/m02400_a3-optimized.cl @@ -25,18 +25,39 @@ void m02400m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke * algorithm specific */ - w[ 4] = 0x80; - w[ 5] = 0; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 16 * 8; - w[15] = 0; + if (pw_len <= 16) + { + w[ 4] = 0x80; + w[ 5] = 0; + w[ 6] = 0; + w[ 7] = 0; + w[ 8] = 0; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 16 * 8; + w[15] = 0; + } + else if (pw_len <= 32) + { + w[ 8] = 0x80; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 32 * 8; + w[15] = 0; + } + else if (pw_len <= 48) + { + w[12] = 0x80; + w[13] = 0; + w[14] = 48 * 8; + w[15] = 0; + } /** * base @@ -219,18 +240,39 @@ void m02400s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke * algorithm specific */ - w[ 4] = 0x80; - w[ 5] = 0; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 16 * 8; - w[15] = 0; + if (pw_len <= 16) + { + w[ 4] = 0x80; + w[ 5] = 0; + w[ 6] = 0; + w[ 7] = 0; + w[ 8] = 0; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 16 * 8; + w[15] = 0; + } + else if (pw_len <= 32) + { + w[ 8] = 0x80; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 32 * 8; + w[15] = 0; + } + else if (pw_len <= 48) + { + w[12] = 0x80; + w[13] = 0; + w[14] = 48 * 8; + w[15] = 0; + } /** * base diff --git a/OpenCL/m02410_a0-optimized.cl b/OpenCL/m02410_a0-optimized.cl index 4db40a0b8..06b77c696 100644 --- a/OpenCL/m02410_a0-optimized.cl +++ b/OpenCL/m02410_a0-optimized.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -70,6 +70,8 @@ __kernel void m02410_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru salt_buf3[2] = 0; salt_buf3[3] = 0; + const u32 salt_len = salt_bufs[salt_pos].salt_len; + /** * loop */ @@ -132,18 +134,41 @@ __kernel void m02410_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru * md5 */ - w1[0] = 0x80; - w1[1] = 0; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 16 * 8; - w3[3] = 0; + const u32x out_salt_len = out_len + salt_len; + + if (out_salt_len <= 16) + { + w1[0] = 0x80; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 16 * 8; + w3[3] = 0; + } + else if (out_salt_len <= 32) + { + w2[0] = 0x80; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 32 * 8; + w3[3] = 0; + } + else if (out_salt_len <= 48) + { + w3[0] = 0x80; + w3[1] = 0; + w3[2] = 48 * 8; + w3[3] = 0; + } u32x a = MD5M_A; u32x b = MD5M_B; @@ -293,6 +318,8 @@ __kernel void m02410_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru salt_buf3[2] = 0; salt_buf3[3] = 0; + const u32 salt_len = salt_bufs[salt_pos].salt_len; + /** * digest */ @@ -367,18 +394,41 @@ __kernel void m02410_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru * md5 */ - w1[0] = 0x80; - w1[1] = 0; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 16 * 8; - w3[3] = 0; + const u32x out_salt_len = out_len + salt_len; + + if (out_salt_len <= 16) + { + w1[0] = 0x80; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 16 * 8; + w3[3] = 0; + } + else if (out_salt_len <= 32) + { + w2[0] = 0x80; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 32 * 8; + w3[3] = 0; + } + else if (out_salt_len <= 48) + { + w3[0] = 0x80; + w3[1] = 0; + w3[2] = 48 * 8; + w3[3] = 0; + } u32x a = MD5M_A; u32x b = MD5M_B; diff --git a/OpenCL/m02410_a1-optimized.cl b/OpenCL/m02410_a1-optimized.cl index a59926a2b..b2cd62be7 100644 --- a/OpenCL/m02410_a1-optimized.cl +++ b/OpenCL/m02410_a1-optimized.cl @@ -3,7 +3,7 @@ * License.....: MIT */ -#define NEW_SIMD_CODE +//#define NEW_SIMD_CODE #include "inc_vendor.cl" #include "inc_hash_constants.h" @@ -68,6 +68,8 @@ __kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rule salt_buf3[2] = 0; salt_buf3[3] = 0; + const u32 salt_len = salt_bufs[salt_pos].salt_len; + /** * loop */ @@ -190,18 +192,41 @@ __kernel void m02410_m04 (__global pw_t *pws, __global const kernel_rule_t *rule * md5 */ - w1[0] = 0x80; - w1[1] = 0; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 16 * 8; - w3[3] = 0; + const u32x pw_salt_len = pw_len + salt_len; + + if (pw_salt_len <= 16) + { + w1[0] = 0x80; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 16 * 8; + w3[3] = 0; + } + else if (pw_salt_len <= 32) + { + w2[0] = 0x80; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 32 * 8; + w3[3] = 0; + } + else if (pw_salt_len <= 48) + { + w3[0] = 0x80; + w3[1] = 0; + w3[2] = 48 * 8; + w3[3] = 0; + } u32x a = MD5M_A; u32x b = MD5M_B; @@ -351,6 +376,8 @@ __kernel void m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rule salt_buf3[2] = 0; salt_buf3[3] = 0; + const u32 salt_len = salt_bufs[salt_pos].salt_len; + /** * digest */ @@ -485,18 +512,41 @@ __kernel void m02410_s04 (__global pw_t *pws, __global const kernel_rule_t *rule * md5 */ - w1[0] = 0x80; - w1[1] = 0; - w1[2] = 0; - w1[3] = 0; - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - w3[0] = 0; - w3[1] = 0; - w3[2] = 16 * 8; - w3[3] = 0; + const u32x pw_salt_len = pw_len + salt_len; + + if (pw_salt_len <= 16) + { + w1[0] = 0x80; + w1[1] = 0; + w1[2] = 0; + w1[3] = 0; + w2[0] = 0; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 16 * 8; + w3[3] = 0; + } + else if (pw_salt_len <= 32) + { + w2[0] = 0x80; + w2[1] = 0; + w2[2] = 0; + w2[3] = 0; + w3[0] = 0; + w3[1] = 0; + w3[2] = 32 * 8; + w3[3] = 0; + } + else if (pw_salt_len <= 48) + { + w3[0] = 0x80; + w3[1] = 0; + w3[2] = 48 * 8; + w3[3] = 0; + } u32x a = MD5M_A; u32x b = MD5M_B; diff --git a/OpenCL/m02410_a3-optimized.cl b/OpenCL/m02410_a3-optimized.cl index fe780074a..e91f4d3fa 100644 --- a/OpenCL/m02410_a3-optimized.cl +++ b/OpenCL/m02410_a3-optimized.cl @@ -47,6 +47,8 @@ void m02410m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke salt_buf3[2] = 0; salt_buf3[3] = 0; + const u32 salt_len = salt_bufs[salt_pos].salt_len; + switch_buffer_by_offset_le_S (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_len); w[ 0] |= salt_buf0[0]; @@ -70,18 +72,41 @@ void m02410m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke * algorithm specific */ - w[ 4] = 0x80; - w[ 5] = 0; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 16 * 8; - w[15] = 0; + const u32 pw_salt_len = pw_len + salt_len; + + if (pw_salt_len <= 16) + { + w[ 4] = 0x80; + w[ 5] = 0; + w[ 6] = 0; + w[ 7] = 0; + w[ 8] = 0; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 16 * 8; + w[15] = 0; + } + else if (pw_salt_len <= 32) + { + w[ 8] = 0x80; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 32 * 8; + w[15] = 0; + } + else if (pw_salt_len <= 48) + { + w[12] = 0x80; + w[13] = 0; + w[14] = 48 * 8; + w[15] = 0; + } /** * base @@ -286,6 +311,8 @@ void m02410s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke salt_buf3[2] = 0; salt_buf3[3] = 0; + const u32 salt_len = salt_bufs[salt_pos].salt_len; + switch_buffer_by_offset_le_S (salt_buf0, salt_buf1, salt_buf2, salt_buf3, pw_len); w[ 0] |= salt_buf0[0]; @@ -309,18 +336,41 @@ void m02410s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global const ke * algorithm specific */ - w[ 4] = 0x80; - w[ 5] = 0; - w[ 6] = 0; - w[ 7] = 0; - w[ 8] = 0; - w[ 9] = 0; - w[10] = 0; - w[11] = 0; - w[12] = 0; - w[13] = 0; - w[14] = 16 * 8; - w[15] = 0; + const u32 pw_salt_len = pw_len + salt_len; + + if (pw_salt_len <= 16) + { + w[ 4] = 0x80; + w[ 5] = 0; + w[ 6] = 0; + w[ 7] = 0; + w[ 8] = 0; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 16 * 8; + w[15] = 0; + } + else if (pw_salt_len <= 32) + { + w[ 8] = 0x80; + w[ 9] = 0; + w[10] = 0; + w[11] = 0; + w[12] = 0; + w[13] = 0; + w[14] = 32 * 8; + w[15] = 0; + } + else if (pw_salt_len <= 48) + { + w[12] = 0x80; + w[13] = 0; + w[14] = 48 * 8; + w[15] = 0; + } /** * base diff --git a/src/interface.c b/src/interface.c index bea9e2664..c73e8628a 100644 --- a/src/interface.c +++ b/src/interface.c @@ -26200,8 +26200,6 @@ int hashconfig_get_pw_max (hashcat_ctx_t *hashcat_ctx, const bool optimized_kern case 112: pw_max = 30; break; // https://www.toadworld.com/platforms/oracle/b/weblog/archive/2013/11/12/oracle-12c-passwords case 1500: pw_max = 8; break; // Underlaying DES max case 2100: pw_max = PW_MAX; break; - case 2400: pw_max = 16; break; // Cisco-PIX MD5 sets w[4] = 0x80 - case 2410: pw_max = 12; break; // Cisco-ASA MD5 sets w[4] = 0x80 plus has a 4 byte fixed salt case 2500: pw_max = 63; break; // WPA/WPA2 limits itself to 63 by RFC case 2501: pw_max = 64; break; // WPA/WPA2 PMK fixed length case 3000: pw_max = 7; break; // LM max diff --git a/tools/test.pl b/tools/test.pl index 2dc3dd281..923ca83c7 100755 --- a/tools/test.pl +++ b/tools/test.pl @@ -45,7 +45,7 @@ use Authen::Passphrase::MySQL323; use Authen::Passphrase::PHPass; use Authen::Passphrase::LANManager; use Encode; -use POSIX qw (strftime); +use POSIX qw (strftime ceil); use Net::DNS::SEC; use Net::DNS::RR::NSEC3; use Convert::EBCDIC qw (ascii2ebcdic); @@ -3516,8 +3516,6 @@ sub passthrough } elsif ($mode == 2410) { - next if length ($word_buf) > 12; - my $salt_len = get_random_num (1, 4); $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, $salt_len)); @@ -4110,7 +4108,7 @@ sub single } elsif ($mode == 2400) { - for (my $i = 1; $i < 16; $i++) + for (my $i = 1; $i < 32; $i++) { if ($len != 0) { @@ -4124,9 +4122,9 @@ sub single } elsif ($mode == 2410) { - my $salt_len = get_random_num (1, 4); + my $salt_len = get_random_num (3, 4); - for (my $i = 1; $i < 13; $i++) + for (my $i = 1; $i < 32; $i++) { if ($len != 0) { @@ -5568,13 +5566,23 @@ sub gen_hash } elsif ($mode == 2400) { - my $hash_buf = Digest::MD5::md5 ($word_buf . "\0" x (16 - length ($word_buf))); + my $word_len = length ($word_buf); + + my $pad_len = ceil ($word_len / 16) * 16; + + my $hash_buf = Digest::MD5::md5 ($word_buf . "\0" x ($pad_len - $word_len)); $tmp_hash = sprintf ("%s", pseudo_base64 ($hash_buf)); } elsif ($mode == 2410) { - my $hash_buf = Digest::MD5::md5 ($word_buf . $salt_buf . "\0" x (16 - length ($word_buf) - length ($salt_buf))); + my $word_salt_buf = $word_buf . $salt_buf; + + my $word_salt_len = length ($word_salt_buf); + + my $pad_len = ceil ($word_salt_len / 16) * 16; + + my $hash_buf = Digest::MD5::md5 ($word_buf . $salt_buf . "\0" x ($pad_len - $word_salt_len)); $tmp_hash = sprintf ("%s:%s", pseudo_base64 ($hash_buf), $salt_buf); } @@ -9569,13 +9577,8 @@ sub rnd my $max = $MAX_LEN; - if ($mode == 2400) + if ($mode == 2410) { - $word_len = min ($word_len, 16); - } - elsif ($mode == 2410) - { - $word_len = min ($word_len, 12); $salt_len = min ($salt_len, 4); }