Init work on Chacha20

OpenCL/inc_types.cl

@@ -779,6 +779,15 @@ typedef struct luks_tmp
 
 } luks_tmp_t;
 
+typedef struct chacha20
+{
+  u32 iv[2];
+  u8  plain[64];
+  u32 plain_length;
+  u32 position;
+
+} chacha20_t;
+
 typedef struct
 {
   int V;

OpenCL/m00670_a0.cl

@@ -0,0 +1,176 @@
+/**
+ * Author......: See docs/credits.txt
+ * License.....: MIT
+ */
+
+#define NEW_SIMD_CODE
+
+#include "inc_vendor.cl"
+#include "inc_hash_constants.h"
+#include "inc_hash_functions.cl"
+#include "inc_types.cl"
+#include "inc_common.cl"
+#include "inc_rp.h"
+#include "inc_rp.cl"
+#include "inc_simd.cl"
+
+__kernel void m00670_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{ 
+  /**
+   * modifier
+   */
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  u32 pw_buf0[4];
+  u32 pw_buf1[4];
+
+  pw_buf0[0] = pws[gid].i[0];
+  pw_buf0[1] = pws[gid].i[1];
+  pw_buf0[2] = pws[gid].i[2];
+  pw_buf0[3] = pws[gid].i[3];
+  pw_buf1[0] = pws[gid].i[4];
+  pw_buf1[1] = pws[gid].i[5];
+  pw_buf1[2] = pws[gid].i[6];
+  pw_buf1[3] = pws[gid].i[7];
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  u32 iv[2];
+
+  iv[0] = esalt_bufs->iv[0];
+  iv[1] = esalt_bufs->iv[1];
+
+  u8 plain[64] = { 0 };
+  u32 plain_length = esalt_bufs->plain_length;
+  u32 position     = esalt_bufs->position;
+
+  for (int i = 0; i < plain_length; i++)
+  {
+    plain[i] = esalt_bufs->plain[i];
+  }
+
+  printf("position: %d, iv: %08x%08x, plain_length: %d, plain: %s\n", position, iv[0], iv[1], plain_length, plain);
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    u32x w0[4] = { 0 };
+    u32x w1[4] = { 0 };
+    u32x w2[4] = { 0 };
+    u32x w3[4] = { 0 };
+
+    const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
+
+    u64x digest[8];
+
+
+    const u32x r0 = h32_from_64(digest[0]);
+    const u32x r1 = l32_from_64(digest[0]);
+    const u32x r2 = h32_from_64(digest[1]);
+    const u32x r3 = l32_from_64(digest[1]);
+
+    COMPARE_M_SIMD(r0, r1, r2, r3);
+  } 
+}
+
+__kernel void m00670_m08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m00670_m16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m00670_s04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{ 
+  /**
+   * modifier
+   */
+
+  const u32 lid = get_local_id (0);
+
+  const u32 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  u32 pw_buf0[4];
+  u32 pw_buf1[4];
+
+  pw_buf0[0] = pws[gid].i[0];
+  pw_buf0[1] = pws[gid].i[1];
+  pw_buf0[2] = pws[gid].i[2];
+  pw_buf0[3] = pws[gid].i[3];
+  pw_buf1[0] = pws[gid].i[4];
+  pw_buf1[1] = pws[gid].i[5];
+  pw_buf1[2] = pws[gid].i[6];
+  pw_buf1[3] = pws[gid].i[7];
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  u32 iv[2];
+
+  iv[0] = esalt_bufs->iv[0];
+  iv[1] = esalt_bufs->iv[1];
+
+  u8 plain[64] = { 0 };
+  u32 plain_length = esalt_bufs->plain_length;
+  u32 position     = esalt_bufs->position;
+
+  for (int i = 0; i < plain_length; i++)
+  {
+    plain[i] = esalt_bufs->plain[i];
+  } 
+
+  printf("s04-> position: %d, iv: %08x%08x, plain_length: %d, plain: %s, cipher: %llu\n", position, iv[0], iv[1], plain_length, plain, digests_buf[digests_offset].digest_buf[0]);
+  
+  /**
+   * digest
+   */
+
+  const u32 search[4] =
+  {
+    digests_buf[digests_offset].digest_buf[DGST_R0],
+    digests_buf[digests_offset].digest_buf[DGST_R1],
+    digests_buf[digests_offset].digest_buf[DGST_R2],
+    digests_buf[digests_offset].digest_buf[DGST_R3]
+  };
+ 
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    u32x w0[4] = { 0 };
+    u32x w1[4] = { 0 };
+    u32x w2[4] = { 0 };
+    u32x w3[4] = { 0 };
+    
+    const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
+
+    u64x digest[8];
+    u64x m[16];
+    u64x v[16];
+
+
+    const u32x r0 = h32_from_64(digest[0]);
+    const u32x r1 = l32_from_64(digest[0]);
+    const u32x r2 = h32_from_64(digest[1]);
+    const u32x r3 = l32_from_64(digest[1]);
+
+    COMPARE_S_SIMD(r0, r1, r2, r3);
+  }  
+}
+
+__kernel void m00670_s08 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}
+
+__kernel void m00670_s16 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const chacha20_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+}

chachaTest.dict

@@ -0,0 +1 @@
+0123456789abcdef0123456789abcdef

chachaTest.hash

@@ -0,0 +1 @@
+$Chacha20$*1*0102030405060708*AAAAAAAAAAA=*E3C0fjxQmkM=

chachaTest.sh

@@ -0,0 +1 @@
+./hashcat -m 670 -a 0 chachaTest.hash chachaTest.dict -u1 -n1 --force --weak-hash-threshold=0 --potfile-disable

include/interface.h

@@ -912,6 +912,8 @@ typedef enum display_len
   DISPLAY_LEN_MAX_501   = 104,
   DISPLAY_LEN_MIN_600   = 8 + 128,
   DISPLAY_LEN_MAX_600   = 8 + 128,
+  DISPLAY_LEN_MIN_670   = 10 + 1 + 1 + 1 + 16 + 1 +  12 + 1 +  12,
+  DISPLAY_LEN_MAX_670   = 10 + 1 + 5 + 1 + 16 + 1 + 128 + 1 + 128,
   DISPLAY_LEN_MIN_900   = 32,
   DISPLAY_LEN_MAX_900   = 32,
   DISPLAY_LEN_MIN_910   = 32 + 1 + 0,
@@ -1329,6 +1331,7 @@ typedef enum hash_type
   HASH_TYPE_ITUNES_BACKUP_10    = 57,
   HASH_TYPE_SKIP32              = 58,
   HASH_TYPE_BLAKE2B             = 59,
+  HASH_TYPE_CHACHA20            = 60,
 
 } hash_type_t;
 
@@ -1353,6 +1356,7 @@ typedef enum kern_type
   KERN_TYPE_PHPASS                  = 400,
   KERN_TYPE_MD5CRYPT                = 500,
   KERN_TYPE_BLAKE2B                 = 600,
+  KERN_TYPE_CHACHA20                = 670,
   KERN_TYPE_MD4                     = 900,
   KERN_TYPE_MD4_PWU                 = 1000,
   KERN_TYPE_MD44_PWUSLT             = 1100,
@@ -1609,6 +1613,7 @@ int postgresql_parse_hash         (u8 *input_buf, u32 input_len, hash_t *hash_bu
 int netscreen_parse_hash          (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
 int keccak_parse_hash             (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
 int blake2b_parse_hash            (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
+int chacha20_parse_hash           (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
 int lm_parse_hash                 (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
 int md4_parse_hash                (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);
 int md5_parse_hash                (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig);

include/types.h

@@ -692,6 +692,15 @@ typedef struct
 
 } blake2_t;
 
+typedef struct
+{
+  u32 iv[2];
+  u8  plain[64];
+  u32 plain_length;
+  u32 position;
+
+} chacha20_t;
+
 typedef struct
 {
   u32 salt_buf[16];

src/interface.c

@@ -96,6 +96,7 @@ static const char HT_00400[] = "phpass, WordPress (MD5), phpBB3 (MD5), Joomla (M
 static const char HT_00500[] = "md5crypt, MD5 (Unix), Cisco-IOS $1$ (MD5)";
 static const char HT_00501[] = "Juniper IVE";
 static const char HT_00600[] = "Blake2-512";
+static const char HT_00670[] = "Chacha20";
 static const char HT_00900[] = "MD4";
 static const char HT_01000[] = "NTLM";
 static const char HT_01100[] = "Domain Cached Credentials (DCC), MS Cache";
@@ -379,6 +380,7 @@ static const char SIGNATURE_FORTIGATE[]        = "AK1";
 static const char SIGNATURE_ATLASSIAN[]        = "{PKCS5S2}";
 static const char SIGNATURE_NETBSD_SHA1CRYPT[] = "$sha1$";
 static const char SIGNATURE_BLAKE2B[]          = "$BLAKE2$";
+static const char SIGNATURE_CHACHA20[]         = "$Chacha20$";
 
 /**
  * decoder / encoder
@@ -5308,6 +5310,49 @@ int blake2b_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UN
   return (PARSER_OK);
 }
 
+int chacha20_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
+{
+  if ((input_len < DISPLAY_LEN_MIN_670) || (input_len > DISPLAY_LEN_MAX_670)) return (PARSER_GLOBAL_LENGTH);
+
+  if (memcmp (SIGNATURE_CHACHA20, input_buf, 10)) return (PARSER_SIGNATURE_UNMATCHED);
+
+  // if (is_valid_hex_string (input_buf + 8, 128) == false) return (PARSER_HASH_ENCODING);
+
+  u8 *digest = (u8 *) hash_buf->digest;
+
+  chacha20_t *chacha20 = (chacha20_t *) hash_buf->esalt;
+
+  u8 *position_marker = (u8 *) strchr ((const char *) input_buf, '*') + 1;
+  if (position_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  u8 *iv_marker = (u8 *) strchr ((const char *) position_marker, '*') + 1;
+  if (iv_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  u8 *plain_marker = (u8 *) strchr ((const char *) iv_marker, '*') + 1;
+  if (plain_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  u8 *cipher_marker = (u8 *) strchr ((const char *) plain_marker, '*') + 1;
+  if (cipher_marker == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  chacha20->position     = 1;
+  chacha20->plain_length = cipher_marker - plain_marker - 1;
+
+  for (int i = 0; i < chacha20->plain_length; i++)
+    chacha20->plain[i] = plain_marker[i];
+
+  chacha20->iv[0] = hex_to_u32 ((const u8 *) iv_marker + 8);
+  chacha20->iv[1] = hex_to_u32 ((const u8 *) iv_marker + 0);
+
+  digest[0] = cipher_marker[ 0];
+  digest[1] = cipher_marker[ 1];
+  digest[2] = cipher_marker[ 2];
+  digest[3] = cipher_marker[ 3];
+
+
+  return (PARSER_OK);
+}
+
+
 int ikepsk_md5_parse_hash (u8 *input_buf, u32 input_len, hash_t *hash_buf, MAYBE_UNUSED const hashconfig_t *hashconfig)
 {
   if ((input_len < DISPLAY_LEN_MIN_5300) || (input_len > DISPLAY_LEN_MAX_5300)) return (PARSER_GLOBAL_LENGTH);
@@ -15039,6 +15084,7 @@ char *strhashtype (const u32 hash_mode)
     case   500: return ((char *) HT_00500);
     case   501: return ((char *) HT_00501);
     case   600: return ((char *) HT_00600);
+    case   670: return ((char *) HT_00670);
     case   900: return ((char *) HT_00900);
     case  1000: return ((char *) HT_01000);
     case  1100: return ((char *) HT_01100);
@@ -18490,6 +18536,29 @@ int ascii_digest (hashcat_ctx_t *hashcat_ctx, char *out_buf, const size_t out_le
         byte_swap_32(ptr[14]),
         byte_swap_32(ptr[15]));
     }
+    else if (hash_type == HASH_TYPE_CHACHA20)
+    {
+      u32 *ptr = digest_buf;
+
+      snprintf (out_buf, out_len - 1, "%s%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x%08x",
+        SIGNATURE_CHACHA20,
+        byte_swap_32(ptr[ 0]),
+        byte_swap_32(ptr[ 1]),
+        byte_swap_32(ptr[ 2]),
+        byte_swap_32(ptr[ 3]),
+        byte_swap_32(ptr[ 4]),
+        byte_swap_32(ptr[ 5]),
+        byte_swap_32(ptr[ 6]),
+        byte_swap_32(ptr[ 7]),
+        byte_swap_32(ptr[ 8]),
+        byte_swap_32(ptr[ 9]),
+        byte_swap_32(ptr[10]),
+        byte_swap_32(ptr[11]),
+        byte_swap_32(ptr[12]),
+        byte_swap_32(ptr[13]),
+        byte_swap_32(ptr[14]),
+        byte_swap_32(ptr[15]));
+    }
     else if (hash_type == HASH_TYPE_RIPEMD160)
     {
       snprintf (out_buf, out_len - 1, "%08x%08x%08x%08x%08x",
@@ -19400,6 +19469,22 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
                  hashconfig->dgst_pos3      = 2;
                  break;
 
+    case   670:  hashconfig->hash_type      = HASH_TYPE_CHACHA20;
+                 hashconfig->salt_type      = SALT_TYPE_EMBEDDED;
+                 hashconfig->attack_exec    = ATTACK_EXEC_INSIDE_KERNEL;
+                 hashconfig->opts_type      = OPTS_TYPE_PT_GENERATE_LE;
+                 hashconfig->kern_type      = KERN_TYPE_CHACHA20;
+                 hashconfig->dgst_size      = DGST_SIZE_8_8;
+                 hashconfig->parse_func     = chacha20_parse_hash;
+                 hashconfig->opti_type      = OPTI_TYPE_ZERO_BYTE
+                                            | OPTI_TYPE_USES_BITS_32
+                                            | OPTI_TYPE_RAW_HASH;
+                 hashconfig->dgst_pos0      = 0;
+                 hashconfig->dgst_pos1      = 1;
+                 hashconfig->dgst_pos2      = 2;
+                 hashconfig->dgst_pos3      = 3;
+                 break;
+
     case   900:  hashconfig->hash_type      = HASH_TYPE_MD4;
                  hashconfig->salt_type      = SALT_TYPE_NONE;
                  hashconfig->attack_exec    = ATTACK_EXEC_INSIDE_KERNEL;
@@ -22687,6 +22772,7 @@ int hashconfig_init (hashcat_ctx_t *hashcat_ctx)
   switch (hashconfig->hash_mode)
   {
     case   600: hashconfig->esalt_size = sizeof (blake2_t);         break;
+    case   670: hashconfig->esalt_size = sizeof (chacha20_t);       break;
     case  2500: hashconfig->esalt_size = sizeof (wpa_t);            break;
     case  5300: hashconfig->esalt_size = sizeof (ikepsk_t);         break;
     case  5400: hashconfig->esalt_size = sizeof (ikepsk_t);         break;