From c941e55a357708468f09b0db5c7512287adfbf25 Mon Sep 17 00:00:00 2001
From: Sein Coray <s3inlc@hashes.org>
Date: Wed, 5 Dec 2018 12:57:54 +0100
Subject: [PATCH] Extended IKE PSK md5/sha1 (-m 5300/5400) to print hashes
 correctly

---
 OpenCL/inc_types.cl           |  2 +-
 OpenCL/m05300_a0-optimized.cl |  8 ++++----
 OpenCL/m05300_a0-pure.cl      |  4 ++--
 OpenCL/m05300_a1-optimized.cl |  8 ++++----
 OpenCL/m05300_a1-pure.cl      |  4 ++--
 OpenCL/m05300_a3-optimized.cl |  8 ++++----
 OpenCL/m05300_a3-pure.cl      |  4 ++--
 OpenCL/m05400_a0-optimized.cl |  8 ++++----
 OpenCL/m05400_a0-pure.cl      |  4 ++--
 OpenCL/m05400_a1-optimized.cl |  8 ++++----
 OpenCL/m05400_a1-pure.cl      |  4 ++--
 OpenCL/m05400_a3-optimized.cl |  8 ++++----
 OpenCL/m05400_a3-pure.cl      |  4 ++--
 include/interface.h           |  2 +-
 src/interface.c               | 30 ++++++++++++++++++++----------
 15 files changed, 58 insertions(+), 48 deletions(-)

diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl
index 4eb5a4c13..8ec01125f 100644
--- a/OpenCL/inc_types.cl
+++ b/OpenCL/inc_types.cl
@@ -1259,7 +1259,7 @@ typedef struct ikepsk
   u32 nr_len;
 
   u32 msg_buf[128];
-  u32 msg_len;
+  u32 msg_len[6];
 
 } ikepsk_t;
 
diff --git a/OpenCL/m05300_a0-optimized.cl b/OpenCL/m05300_a0-optimized.cl
index 40d80eb99..c50749508 100644
--- a/OpenCL/m05300_a0-optimized.cl
+++ b/OpenCL/m05300_a0-optimized.cl
@@ -155,7 +155,7 @@ __kernel void m05300_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * loop
@@ -222,7 +222,7 @@ __kernel void m05300_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0[0] = s_msg_buf[off +  0];
       w0[1] = s_msg_buf[off +  1];
@@ -330,7 +330,7 @@ __kernel void m05300_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * digest
@@ -409,7 +409,7 @@ __kernel void m05300_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0[0] = s_msg_buf[off +  0];
       w0[1] = s_msg_buf[off +  1];
diff --git a/OpenCL/m05300_a0-pure.cl b/OpenCL/m05300_a0-pure.cl
index d336edb91..51e56415a 100644
--- a/OpenCL/m05300_a0-pure.cl
+++ b/OpenCL/m05300_a0-pure.cl
@@ -76,7 +76,7 @@ __kernel void m05300_mxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
 
     md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     md5_hmac_final (&ctx);
 
@@ -162,7 +162,7 @@ __kernel void m05300_sxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
 
     md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     md5_hmac_final (&ctx);
 
diff --git a/OpenCL/m05300_a1-optimized.cl b/OpenCL/m05300_a1-optimized.cl
index 1738a951a..5fa7542a0 100644
--- a/OpenCL/m05300_a1-optimized.cl
+++ b/OpenCL/m05300_a1-optimized.cl
@@ -153,7 +153,7 @@ __kernel void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * loop
@@ -280,7 +280,7 @@ __kernel void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t))
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0[0] = s_msg_buf[off +  0];
       w0[1] = s_msg_buf[off +  1];
@@ -388,7 +388,7 @@ __kernel void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * digest
@@ -527,7 +527,7 @@ __kernel void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t))
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0[0] = s_msg_buf[off +  0];
       w0[1] = s_msg_buf[off +  1];
diff --git a/OpenCL/m05300_a1-pure.cl b/OpenCL/m05300_a1-pure.cl
index 6e4453d8b..256e397ee 100644
--- a/OpenCL/m05300_a1-pure.cl
+++ b/OpenCL/m05300_a1-pure.cl
@@ -99,7 +99,7 @@ __kernel void m05300_mxx (KERN_ATTR_ESALT (ikepsk_t))
 
     md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     md5_hmac_final (&ctx);
 
@@ -210,7 +210,7 @@ __kernel void m05300_sxx (KERN_ATTR_ESALT (ikepsk_t))
 
     md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     md5_hmac_final (&ctx);
 
diff --git a/OpenCL/m05300_a3-optimized.cl b/OpenCL/m05300_a3-optimized.cl
index 99fb7d6ac..702c28775 100644
--- a/OpenCL/m05300_a3-optimized.cl
+++ b/OpenCL/m05300_a3-optimized.cl
@@ -112,7 +112,7 @@ DECLSPEC void m05300m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * loop
@@ -200,7 +200,7 @@ DECLSPEC void m05300m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0_t[0] = s_msg_buf[off +  0];
       w0_t[1] = s_msg_buf[off +  1];
@@ -259,7 +259,7 @@ DECLSPEC void m05300s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * digest
@@ -359,7 +359,7 @@ DECLSPEC void m05300s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0_t[0] = s_msg_buf[off +  0];
       w0_t[1] = s_msg_buf[off +  1];
diff --git a/OpenCL/m05300_a3-pure.cl b/OpenCL/m05300_a3-pure.cl
index c2e5384ca..27ed51150 100644
--- a/OpenCL/m05300_a3-pure.cl
+++ b/OpenCL/m05300_a3-pure.cl
@@ -85,7 +85,7 @@ __kernel void m05300_mxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
 
     md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     md5_hmac_final (&ctx);
 
@@ -182,7 +182,7 @@ __kernel void m05300_sxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
 
     md5_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    md5_hmac_update_global (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     md5_hmac_final (&ctx);
 
diff --git a/OpenCL/m05400_a0-optimized.cl b/OpenCL/m05400_a0-optimized.cl
index 64240619d..aeecadfbb 100644
--- a/OpenCL/m05400_a0-optimized.cl
+++ b/OpenCL/m05400_a0-optimized.cl
@@ -159,7 +159,7 @@ __kernel void m05400_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * loop
@@ -235,7 +235,7 @@ __kernel void m05400_m04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0[0] = s_msg_buf[off +  0];
       w0[1] = s_msg_buf[off +  1];
@@ -343,7 +343,7 @@ __kernel void m05400_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * digest
@@ -431,7 +431,7 @@ __kernel void m05400_s04 (KERN_ATTR_RULES_ESALT (ikepsk_t))
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0[0] = s_msg_buf[off +  0];
       w0[1] = s_msg_buf[off +  1];
diff --git a/OpenCL/m05400_a0-pure.cl b/OpenCL/m05400_a0-pure.cl
index 7c14d9c46..8a0e1f80b 100644
--- a/OpenCL/m05400_a0-pure.cl
+++ b/OpenCL/m05400_a0-pure.cl
@@ -76,7 +76,7 @@ __kernel void m05400_mxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
 
     sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     sha1_hmac_final (&ctx);
 
@@ -162,7 +162,7 @@ __kernel void m05400_sxx (KERN_ATTR_RULES_ESALT (ikepsk_t))
 
     sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     sha1_hmac_final (&ctx);
 
diff --git a/OpenCL/m05400_a1-optimized.cl b/OpenCL/m05400_a1-optimized.cl
index 9d2bcf127..9877e0768 100644
--- a/OpenCL/m05400_a1-optimized.cl
+++ b/OpenCL/m05400_a1-optimized.cl
@@ -157,7 +157,7 @@ __kernel void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * loop
@@ -301,7 +301,7 @@ __kernel void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t))
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0[0] = s_msg_buf[off +  0];
       w0[1] = s_msg_buf[off +  1];
@@ -409,7 +409,7 @@ __kernel void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * digest
@@ -565,7 +565,7 @@ __kernel void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t))
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0[0] = s_msg_buf[off +  0];
       w0[1] = s_msg_buf[off +  1];
diff --git a/OpenCL/m05400_a1-pure.cl b/OpenCL/m05400_a1-pure.cl
index a79e5a219..ae3cfe10b 100644
--- a/OpenCL/m05400_a1-pure.cl
+++ b/OpenCL/m05400_a1-pure.cl
@@ -99,7 +99,7 @@ __kernel void m05400_mxx (KERN_ATTR_ESALT (ikepsk_t))
 
     sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     sha1_hmac_final (&ctx);
 
@@ -210,7 +210,7 @@ __kernel void m05400_sxx (KERN_ATTR_ESALT (ikepsk_t))
 
     sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     sha1_hmac_final (&ctx);
 
diff --git a/OpenCL/m05400_a3-optimized.cl b/OpenCL/m05400_a3-optimized.cl
index 53e346845..321ab745c 100644
--- a/OpenCL/m05400_a3-optimized.cl
+++ b/OpenCL/m05400_a3-optimized.cl
@@ -116,7 +116,7 @@ DECLSPEC void m05400m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * loop
@@ -204,7 +204,7 @@ DECLSPEC void m05400m (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0_t[0] = s_msg_buf[off +  0];
       w0_t[1] = s_msg_buf[off +  1];
@@ -263,7 +263,7 @@ DECLSPEC void m05400s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
    */
 
   const u32 nr_len  = esalt_bufs[digests_offset].nr_len;
-  const u32 msg_len = esalt_bufs[digests_offset].msg_len;
+  const u32 msg_len = esalt_bufs[digests_offset].msg_len[5];
 
   /**
    * digest
@@ -363,7 +363,7 @@ DECLSPEC void m05400s (u32 *w0, u32 *w1, u32 *w2, u32 *w3, const u32 pw_len, KER
     int left;
     int off;
 
-    for (left = esalt_bufs[digests_offset].msg_len, off = 0; left >= 56; left -= 64, off += 16)
+    for (left = msg_len, off = 0; left >= 56; left -= 64, off += 16)
     {
       w0_t[0] = s_msg_buf[off +  0];
       w0_t[1] = s_msg_buf[off +  1];
diff --git a/OpenCL/m05400_a3-pure.cl b/OpenCL/m05400_a3-pure.cl
index e7e303036..1c9237e2c 100644
--- a/OpenCL/m05400_a3-pure.cl
+++ b/OpenCL/m05400_a3-pure.cl
@@ -85,7 +85,7 @@ __kernel void m05400_mxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
 
     sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     sha1_hmac_final (&ctx);
 
@@ -182,7 +182,7 @@ __kernel void m05400_sxx (KERN_ATTR_VECTOR_ESALT (ikepsk_t))
 
     sha1_hmac_init_64 (&ctx, w0, w1, w2, w3);
 
-    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len);
+    sha1_hmac_update_global_swap (&ctx, esalt_bufs[digests_offset].msg_buf, esalt_bufs[digests_offset].msg_len[5]);
 
     sha1_hmac_final (&ctx);
 
diff --git a/include/interface.h b/include/interface.h
index 4074ef995..0685ede9f 100644
--- a/include/interface.h
+++ b/include/interface.h
@@ -250,7 +250,7 @@ typedef struct ikepsk
   u32 nr_len;
 
   u32 msg_buf[128];
-  u32 msg_len;
+  u32 msg_len[6];
 
 } ikepsk_t;
 
diff --git a/src/interface.c b/src/interface.c
index c96719a0b..4eec6b01f 100644
--- a/src/interface.c
+++ b/src/interface.c
@@ -6705,10 +6705,15 @@ int ikepsk_md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE
 
   if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
 
-  ikepsk->msg_len = (token.len[0] + token.len[1] + token.len[2] + token.len[3] + token.len[4] + token.len[5]) / 2;
+  ikepsk->msg_len[0] =                      token.len[0] / 2;
+  ikepsk->msg_len[1] = ikepsk->msg_len[0] + token.len[1] / 2;
+  ikepsk->msg_len[2] = ikepsk->msg_len[1] + token.len[2] / 2;
+  ikepsk->msg_len[3] = ikepsk->msg_len[2] + token.len[3] / 2;
+  ikepsk->msg_len[4] = ikepsk->msg_len[3] + token.len[4] / 2;
+  ikepsk->msg_len[5] = ikepsk->msg_len[4] + token.len[5] / 2;
   ikepsk->nr_len  = (token.len[6] + token.len[7]) / 2;
 
-  if (ikepsk->msg_len > 512) return (PARSER_SALT_LENGTH);
+  if (ikepsk->msg_len[5] > 512) return (PARSER_SALT_LENGTH);
   if (ikepsk->nr_len  > 64)  return (PARSER_SALT_LENGTH);
 
   u8 *ptr1 = (u8 *) ikepsk->msg_buf;
@@ -6821,10 +6826,15 @@ int ikepsk_sha1_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYB
 
   if (rc_tokenizer != PARSER_OK) return (rc_tokenizer);
 
-  ikepsk->msg_len = (token.len[0] + token.len[1] + token.len[2] + token.len[3] + token.len[4] + token.len[5]) / 2;
+  ikepsk->msg_len[0] =                      token.len[0] / 2;
+  ikepsk->msg_len[1] = ikepsk->msg_len[0] + token.len[1] / 2;
+  ikepsk->msg_len[2] = ikepsk->msg_len[1] + token.len[2] / 2;
+  ikepsk->msg_len[3] = ikepsk->msg_len[2] + token.len[3] / 2;
+  ikepsk->msg_len[4] = ikepsk->msg_len[3] + token.len[4] / 2;
+  ikepsk->msg_len[5] = ikepsk->msg_len[4] + token.len[5] / 2;
   ikepsk->nr_len  = (token.len[6] + token.len[7]) / 2;
 
-  if (ikepsk->msg_len > 512) return (PARSER_SALT_LENGTH);
+  if (ikepsk->msg_len[5] > 512) return (PARSER_SALT_LENGTH);
   if (ikepsk->nr_len  > 64)  return (PARSER_SALT_LENGTH);
 
   u8 *ptr1 = (u8 *) ikepsk->msg_buf;
@@ -20037,11 +20047,11 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
 
     // msg_buf
 
-    u32 ikepsk_msg_len = ikepsk->msg_len / 4;
+    u32 ikepsk_msg_len = ikepsk->msg_len[5] / 4;
 
     for (u32 i = 0; i < ikepsk_msg_len; i++)
     {
-      if ((i == 32) || (i == 64) || (i == 66) || (i == 68) || (i == 108))
+      if ((i == ikepsk->msg_len[0] / 4) || (i == ikepsk->msg_len[1] / 4) || (i == ikepsk->msg_len[2] / 4) || (i == ikepsk->msg_len[3] / 4) || (i == ikepsk->msg_len[4] / 4))
       {
         snprintf (out_buf, buf_len, ":");
 
@@ -20103,11 +20113,11 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
 
     // msg_buf
 
-    u32 ikepsk_msg_len = ikepsk->msg_len / 4;
+    u32 ikepsk_msg_len = ikepsk->msg_len[5] / 4;
 
     for (u32 i = 0; i < ikepsk_msg_len; i++)
     {
-      if ((i == 32) || (i == 64) || (i == 66) || (i == 68) || (i == 108))
+      if ((i == ikepsk->msg_len[0] / 4) || (i == ikepsk->msg_len[1] / 4) || (i == ikepsk->msg_len[2] / 4) || (i == ikepsk->msg_len[3] / 4) || (i == ikepsk->msg_len[4] / 4))
       {
         snprintf (out_buf, buf_len, ":");
 
@@ -29509,10 +29519,10 @@ void hashconfig_benchmark_defaults (hashcat_ctx_t *hashcat_ctx, salt_t *salt, vo
       case  2501: ((wpa_eapol_t *)        esalt)->eapol_len     = 128;
                   break;
       case  5300: ((ikepsk_t *)           esalt)->nr_len        = 1;
-                  ((ikepsk_t *)           esalt)->msg_len       = 1;
+                  ((ikepsk_t *)           esalt)->msg_len[5]    = 1;
                   break;
       case  5400: ((ikepsk_t *)           esalt)->nr_len        = 1;
-                  ((ikepsk_t *)           esalt)->msg_len       = 1;
+                  ((ikepsk_t *)           esalt)->msg_len[5]    = 1;
                   break;
       case  5500: ((netntlm_t *)          esalt)->user_len      = 1;
                   ((netntlm_t *)          esalt)->domain_len    = 1;