From 152f0b5152883051c232339256f7f84c87eb6a07 Mon Sep 17 00:00:00 2001 From: DoZ10 Date: Thu, 4 May 2017 22:34:52 -0400 Subject: [PATCH] Init work on Chacha20 --- OpenCL/inc_types.cl | 9 +++ OpenCL/m00670_a0.cl | 176 ++++++++++++++++++++++++++++++++++++++++++++ chachaTest.dict | 1 + chachaTest.hash | 1 + chachaTest.sh | 1 + include/interface.h | 5 ++ include/types.h | 9 +++ src/interface.c | 86 ++++++++++++++++++++++ 8 files changed, 288 insertions(+) create mode 100644 OpenCL/m00670_a0.cl create mode 100644 chachaTest.dict create mode 100644 chachaTest.hash create mode 100755 chachaTest.sh diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index c4d69dbe7..b5d7daad6 100644 --- a/OpenCL/inc_types.cl +++ b/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; diff --git a/OpenCL/m00670_a0.cl b/OpenCL/m00670_a0.cl new file mode 100644 index 000000000..2665a32c5 --- /dev/null +++ b/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) +{ +} diff --git a/chachaTest.dict b/chachaTest.dict new file mode 100644 index 000000000..cc6718945 --- /dev/null +++ b/chachaTest.dict @@ -0,0 +1 @@ +0123456789abcdef0123456789abcdef diff --git a/chachaTest.hash b/chachaTest.hash new file mode 100644 index 000000000..04c72e265 --- /dev/null +++ b/chachaTest.hash @@ -0,0 +1 @@ +$Chacha20$*1*0102030405060708*AAAAAAAAAAA=*E3C0fjxQmkM= diff --git a/chachaTest.sh b/chachaTest.sh new file mode 100755 index 000000000..253f5f357 --- /dev/null +++ b/chachaTest.sh @@ -0,0 +1 @@ +./hashcat -m 670 -a 0 chachaTest.hash chachaTest.dict -u1 -n1 --force --weak-hash-threshold=0 --potfile-disable diff --git a/include/interface.h b/include/interface.h index e9c6fbc2a..25c7657bf 100644 --- a/include/interface.h +++ b/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); diff --git a/include/types.h b/include/types.h index 570073fc0..93b2a0ec8 100644 --- a/include/types.h +++ b/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]; diff --git a/src/interface.c b/src/interface.c index 5b040df89..0068b38e5 100644 --- a/src/interface.c +++ b/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;