github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-12-27 05:13:45 +01:00
Code Releases Activity

OpenCL Kernels: Remove password length restriction to 16 for Cisco-PIX and Cisco-ASA hashes

Browse Source 
Fixes #1488
This commit is contained in:
jsteube 2018-01-27 22:21:44 +01:00
parent 4c981ea0ba
commit e877c30ebc
8 changed files with 441 additions and 164 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

92
OpenCL/m02400_a0-optimized.cl
View File

@ -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;

92
OpenCL/m02400_a1-optimized.cl
View File

@ -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;

90
OpenCL/m02400_a3-optimized.cl
View File

@ -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

100
OpenCL/m02410_a0-optimized.cl
View File

@ -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;

100
OpenCL/m02410_a1-optimized.cl
View File

@ -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;

98
OpenCL/m02410_a3-optimized.cl
View File

@ -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

2
src/interface.c
View File

@ -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

31
tools/test.pl
View File

@ -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);
}