fixes #1624: increase esalt/nonce buffer to 1024 for -m 11400 = SIP

OpenCL/inc_types.cl

@@ -1234,7 +1234,7 @@ typedef struct sip
   u32 salt_buf[32];
   u32 salt_len;
 
-  u32 esalt_buf[48];
+  u32 esalt_buf[256];
   u32 esalt_len;
 
 } sip_t;

OpenCL/m11400_a3-pure.cl

@@ -69,7 +69,7 @@ __kernel void m11400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 esalt_len = esalt_bufs[digests_offset].esalt_len;
 
-  u32x esalt_buf[48] = { 0 };
+  u32x esalt_buf[256] = { 0 };
 
   for (int i = 0, idx = 0; i < esalt_len; i += 4, idx += 1)
   {
@@ -201,7 +201,7 @@ __kernel void m11400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 esalt_len = esalt_bufs[digests_offset].esalt_len;
 
-  u32x esalt_buf[48] = { 0 };
+  u32x esalt_buf[256] = { 0 };
 
   for (int i = 0, idx = 0; i < esalt_len; i += 4, idx += 1)
   {

docs/changes.txt

@@ -17,6 +17,7 @@
 - OpenCL kernels: Add '-pure' prefix to kernel filenames to avoid problems caused by reusing existing hashcat installation folder
 - OpenCL kernels: Removed the use of 'volatile' in inline assembly instructions where it is not needed
 - OpenCL kernels: Switched array pointer types in function declarations in order to be compatible with OpenCL 2.0
+- SIP cracking: Increased the nonce field to allow a salt of 1024 bytes
 
 ##
 ## Bugs

include/interface.h

@@ -216,7 +216,7 @@ typedef struct sip
   u32 salt_buf[32];
   u32 salt_len;
 
-  u32 esalt_buf[48];
+  u32 esalt_buf[256];
   u32 esalt_len;
 
 } sip_t;

src/interface.c

@@ -12576,25 +12576,25 @@ int sip_auth_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_U
 
   token.sep[9]      = '*';
   token.len_min[9]  = 1;
-  token.len_max[9]  = 50;
+  token.len_max[9]  = 1024;
   token.attr[9]     = TOKEN_ATTR_VERIFY_LENGTH
                     | TOKEN_ATTR_TERMINATE_STRING;
 
   token.sep[10]     = '*';
   token.len_min[10] = 0;
-  token.len_max[10] = 50;
+  token.len_max[10] = 1024;
   token.attr[10]    = TOKEN_ATTR_VERIFY_LENGTH
                     | TOKEN_ATTR_TERMINATE_STRING;
 
   token.sep[11]     = '*';
   token.len_min[11] = 0;
-  token.len_max[11] = 50;
+  token.len_max[11] = 1024;
   token.attr[11]    = TOKEN_ATTR_VERIFY_LENGTH
                     | TOKEN_ATTR_TERMINATE_STRING;
 
   token.sep[12]     = '*';
   token.len_min[12] = 0;
-  token.len_max[12] = 50;
+  token.len_max[12] = 1024;
   token.attr[12]    = TOKEN_ATTR_VERIFY_LENGTH
                     | TOKEN_ATTR_TERMINATE_STRING;
 
@@ -12711,7 +12711,7 @@ int sip_auth_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_U
 
   u32 esalt_len = 0;
 
-  u32 max_esalt_len = sizeof (sip->esalt_buf); // 151 = (64 + 64 + 55) - 32, where 32 is the hexadecimal MD5 HA1 hash
+  u32 max_esalt_len = sizeof (sip->esalt_buf);
 
   // there are 2 possibilities for the esalt:
 
@@ -12753,7 +12753,7 @@ int sip_auth_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_U
   {
     esalt_len = 1 + nonce_len + 1 + 32;
 
-    //if (esalt_len > max_esalt_len) return (PARSER_SALT_LENGTH);
+    if (esalt_len > max_esalt_len) return (PARSER_SALT_LENGTH);
 
     snprintf ((char *) esalt_buf_ptr, max_esalt_len, ":%s:%08x%08x%08x%08x",
       nonce_pos,
@@ -12763,8 +12763,6 @@ int sip_auth_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_U
       tmp_digest[3]);
   }
 
-  if (esalt_len >= 152) return (PARSER_SALT_LENGTH);
-
   // add 0x80 to esalt
 
   esalt_buf_ptr[esalt_len] = 0x80;