Added new hash mode -m 13000 = RAR5

diff --git a/OpenCL/m13000.cl b/OpenCL/m13000.cl
new file mode 100644 (file)
index 0000000..383ade1
--- /dev/null
+++ b/OpenCL/m13000.cl
@@ -0,0 +1,528 @@
+/**
+ * Author......: Jens Steube <jens.steube@gmail.com>
+ * License.....: MIT
+ */
+
+#define _RAR5_
+
+#include "include/constants.h"
+#include "include/kernel_vendor.h"
+
+#define DGST_R0 0
+#define DGST_R1 1
+#define DGST_R2 2
+#define DGST_R3 3
+
+#include "include/kernel_functions.c"
+
+#include "OpenCL/types_ocl.c"
+#include "OpenCL/common.c"
+
+#define COMPARE_S "OpenCL/check_single_comp4.c"
+#define COMPARE_M "OpenCL/check_multi_comp4.c"
+
+__constant u32 k_sha256[64] =
+{
+  SHA256C00, SHA256C01, SHA256C02, SHA256C03,
+  SHA256C04, SHA256C05, SHA256C06, SHA256C07,
+  SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
+  SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
+  SHA256C10, SHA256C11, SHA256C12, SHA256C13,
+  SHA256C14, SHA256C15, SHA256C16, SHA256C17,
+  SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
+  SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
+  SHA256C20, SHA256C21, SHA256C22, SHA256C23,
+  SHA256C24, SHA256C25, SHA256C26, SHA256C27,
+  SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
+  SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
+  SHA256C30, SHA256C31, SHA256C32, SHA256C33,
+  SHA256C34, SHA256C35, SHA256C36, SHA256C37,
+  SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
+  SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
+};
+
+static void sha256_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[8])
+{
+  u32 a = digest[0];
+  u32 b = digest[1];
+  u32 c = digest[2];
+  u32 d = digest[3];
+  u32 e = digest[4];
+  u32 f = digest[5];
+  u32 g = digest[6];
+  u32 h = digest[7];
+
+  u32 w0_t = w0[0];
+  u32 w1_t = w0[1];
+  u32 w2_t = w0[2];
+  u32 w3_t = w0[3];
+  u32 w4_t = w1[0];
+  u32 w5_t = w1[1];
+  u32 w6_t = w1[2];
+  u32 w7_t = w1[3];
+  u32 w8_t = w2[0];
+  u32 w9_t = w2[1];
+  u32 wa_t = w2[2];
+  u32 wb_t = w2[3];
+  u32 wc_t = w3[0];
+  u32 wd_t = w3[1];
+  u32 we_t = w3[2];
+  u32 wf_t = w3[3];
+
+  #define ROUND_EXPAND()                            \
+  {                                                 \
+    w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t);  \
+    w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t);  \
+    w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t);  \
+    w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t);  \
+    w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t);  \
+    w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t);  \
+    w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t);  \
+    w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t);  \
+    w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t);  \
+    w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t);  \
+    wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t);  \
+    wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t);  \
+    wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t);  \
+    wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t);  \
+    we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t);  \
+    wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t);  \
+  }
+
+  #define ROUND_STEP(i)                                                                   \
+  {                                                                                       \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i +  0]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i +  1]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i +  2]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i +  3]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i +  4]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i +  5]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i +  6]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i +  7]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i +  8]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i +  9]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
+    SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
+  }
+
+  ROUND_STEP (0);
+
+  #pragma unroll
+  for (int i = 16; i < 64; i += 16)
+  {
+    ROUND_EXPAND (); ROUND_STEP (i);
+  }
+
+  digest[0] += a;
+  digest[1] += b;
+  digest[2] += c;
+  digest[3] += d;
+  digest[4] += e;
+  digest[5] += f;
+  digest[6] += g;
+  digest[7] += h;
+}
+
+static void hmac_sha256_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u32 opad[8])
+{
+  w0[0] = w0[0] ^ 0x36363636;
+  w0[1] = w0[1] ^ 0x36363636;
+  w0[2] = w0[2] ^ 0x36363636;
+  w0[3] = w0[3] ^ 0x36363636;
+  w1[0] = w1[0] ^ 0x36363636;
+  w1[1] = w1[1] ^ 0x36363636;
+  w1[2] = w1[2] ^ 0x36363636;
+  w1[3] = w1[3] ^ 0x36363636;
+  w2[0] = w2[0] ^ 0x36363636;
+  w2[1] = w2[1] ^ 0x36363636;
+  w2[2] = w2[2] ^ 0x36363636;
+  w2[3] = w2[3] ^ 0x36363636;
+  w3[0] = w3[0] ^ 0x36363636;
+  w3[1] = w3[1] ^ 0x36363636;
+  w3[2] = w3[2] ^ 0x36363636;
+  w3[3] = w3[3] ^ 0x36363636;
+
+  ipad[0] = SHA256M_A;
+  ipad[1] = SHA256M_B;
+  ipad[2] = SHA256M_C;
+  ipad[3] = SHA256M_D;
+  ipad[4] = SHA256M_E;
+  ipad[5] = SHA256M_F;
+  ipad[6] = SHA256M_G;
+  ipad[7] = SHA256M_H;
+
+  sha256_transform (w0, w1, w2, w3, ipad);
+
+  w0[0] = w0[0] ^ 0x6a6a6a6a;
+  w0[1] = w0[1] ^ 0x6a6a6a6a;
+  w0[2] = w0[2] ^ 0x6a6a6a6a;
+  w0[3] = w0[3] ^ 0x6a6a6a6a;
+  w1[0] = w1[0] ^ 0x6a6a6a6a;
+  w1[1] = w1[1] ^ 0x6a6a6a6a;
+  w1[2] = w1[2] ^ 0x6a6a6a6a;
+  w1[3] = w1[3] ^ 0x6a6a6a6a;
+  w2[0] = w2[0] ^ 0x6a6a6a6a;
+  w2[1] = w2[1] ^ 0x6a6a6a6a;
+  w2[2] = w2[2] ^ 0x6a6a6a6a;
+  w2[3] = w2[3] ^ 0x6a6a6a6a;
+  w3[0] = w3[0] ^ 0x6a6a6a6a;
+  w3[1] = w3[1] ^ 0x6a6a6a6a;
+  w3[2] = w3[2] ^ 0x6a6a6a6a;
+  w3[3] = w3[3] ^ 0x6a6a6a6a;
+
+  opad[0] = SHA256M_A;
+  opad[1] = SHA256M_B;
+  opad[2] = SHA256M_C;
+  opad[3] = SHA256M_D;
+  opad[4] = SHA256M_E;
+  opad[5] = SHA256M_F;
+  opad[6] = SHA256M_G;
+  opad[7] = SHA256M_H;
+
+  sha256_transform (w0, w1, w2, w3, opad);
+}
+
+static void hmac_sha256_run (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[8], u32 opad[8], u32 digest[8])
+{
+  digest[0] = ipad[0];
+  digest[1] = ipad[1];
+  digest[2] = ipad[2];
+  digest[3] = ipad[3];
+  digest[4] = ipad[4];
+  digest[5] = ipad[5];
+  digest[6] = ipad[6];
+  digest[7] = ipad[7];
+
+  sha256_transform (w0, w1, w2, w3, digest);
+
+  w0[0] = digest[0];
+  w0[1] = digest[1];
+  w0[2] = digest[2];
+  w0[3] = digest[3];
+  w1[0] = digest[4];
+  w1[1] = digest[5];
+  w1[2] = digest[6];
+  w1[3] = digest[7];
+  w2[0] = 0x80000000;
+  w2[1] = 0;
+  w2[2] = 0;
+  w2[3] = 0;
+  w3[0] = 0;
+  w3[1] = 0;
+  w3[2] = 0;
+  w3[3] = (64 + 32) * 8;
+
+  digest[0] = opad[0];
+  digest[1] = opad[1];
+  digest[2] = opad[2];
+  digest[3] = opad[3];
+  digest[4] = opad[4];
+  digest[5] = opad[5];
+  digest[6] = opad[6];
+  digest[7] = opad[7];
+
+  sha256_transform (w0, w1, w2, w3, digest);
+}
+
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m13000_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha256_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global pbkdf2_sha256_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * base
+   */
+
+  const u32 gid = get_global_id (0);
+  const u32 lid = get_local_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  u32 w0[4];
+
+  w0[0] = pws[gid].i[ 0];
+  w0[1] = pws[gid].i[ 1];
+  w0[2] = pws[gid].i[ 2];
+  w0[3] = pws[gid].i[ 3];
+
+  u32 w1[4];
+
+  w1[0] = pws[gid].i[ 4];
+  w1[1] = pws[gid].i[ 5];
+  w1[2] = pws[gid].i[ 6];
+  w1[3] = pws[gid].i[ 7];
+
+  u32 w2[4];
+
+  w2[0] = pws[gid].i[ 8];
+  w2[1] = pws[gid].i[ 9];
+  w2[2] = pws[gid].i[10];
+  w2[3] = pws[gid].i[11];
+
+  u32 w3[4];
+
+  w3[0] = pws[gid].i[12];
+  w3[1] = pws[gid].i[13];
+  w3[2] = pws[gid].i[14];
+  w3[3] = pws[gid].i[15];
+
+  const u32 pw_len = pws[gid].pw_len;
+
+  w0[0] = swap32 (w0[0]);
+  w0[1] = swap32 (w0[1]);
+  w0[2] = swap32 (w0[2]);
+  w0[3] = swap32 (w0[3]);
+  w1[0] = swap32 (w1[0]);
+  w1[1] = swap32 (w1[1]);
+  w1[2] = swap32 (w1[2]);
+  w1[3] = swap32 (w1[3]);
+  w2[0] = swap32 (w2[0]);
+  w2[1] = swap32 (w2[1]);
+  w2[2] = swap32 (w2[2]);
+  w2[3] = swap32 (w2[3]);
+  w3[0] = swap32 (w3[0]);
+  w3[1] = swap32 (w3[1]);
+  w3[2] = swap32 (w3[2]);
+  w3[3] = swap32 (w3[3]);
+
+  /**
+   * salt
+   */
+
+  u32 salt_len = 16;
+
+  u32 salt_buf0[4];
+
+  salt_buf0[0] = salt_bufs[salt_pos].salt_buf[0];
+  salt_buf0[1] = salt_bufs[salt_pos].salt_buf[1];
+  salt_buf0[2] = salt_bufs[salt_pos].salt_buf[2];
+  salt_buf0[3] = salt_bufs[salt_pos].salt_buf[3];
+
+  u32 salt_buf1[4];
+
+  salt_buf1[0] = 0x00000001;
+  salt_buf1[1] = 0x80000000;
+  salt_buf1[2] = 0;
+  salt_buf1[3] = 0;
+
+  u32 salt_buf2[4];
+
+  salt_buf2[0] = 0;
+  salt_buf2[1] = 0;
+  salt_buf2[2] = 0;
+  salt_buf2[3] = 0;
+
+  u32 salt_buf3[4];
+
+  salt_buf3[0] = 0;
+  salt_buf3[1] = 0;
+  salt_buf3[2] = 0;
+  salt_buf3[3] = (64 + salt_len + 4) * 8;
+
+  u32 ipad[8];
+  u32 opad[8];
+
+  hmac_sha256_pad (w0, w1, w2, w3, ipad, opad);
+
+  tmps[gid].ipad[0] = ipad[0];
+  tmps[gid].ipad[1] = ipad[1];
+  tmps[gid].ipad[2] = ipad[2];
+  tmps[gid].ipad[3] = ipad[3];
+  tmps[gid].ipad[4] = ipad[4];
+  tmps[gid].ipad[5] = ipad[5];
+  tmps[gid].ipad[6] = ipad[6];
+  tmps[gid].ipad[7] = ipad[7];
+
+  tmps[gid].opad[0] = opad[0];
+  tmps[gid].opad[1] = opad[1];
+  tmps[gid].opad[2] = opad[2];
+  tmps[gid].opad[3] = opad[3];
+  tmps[gid].opad[4] = opad[4];
+  tmps[gid].opad[5] = opad[5];
+  tmps[gid].opad[6] = opad[6];
+  tmps[gid].opad[7] = opad[7];
+
+  for (u32 i = 0, j = 1; i < 8; i += 8, j += 1)
+  {
+    u32 dgst[8];
+
+    hmac_sha256_run (salt_buf0, salt_buf1, salt_buf2, salt_buf3, ipad, opad, dgst);
+
+    tmps[gid].dgst[i + 0] = dgst[0];
+    tmps[gid].dgst[i + 1] = dgst[1];
+    tmps[gid].dgst[i + 2] = dgst[2];
+    tmps[gid].dgst[i + 3] = dgst[3];
+    tmps[gid].dgst[i + 4] = dgst[4];
+    tmps[gid].dgst[i + 5] = dgst[5];
+    tmps[gid].dgst[i + 6] = dgst[6];
+    tmps[gid].dgst[i + 7] = dgst[7];
+
+    tmps[gid].out[i + 0] = dgst[0];
+    tmps[gid].out[i + 1] = dgst[1];
+    tmps[gid].out[i + 2] = dgst[2];
+    tmps[gid].out[i + 3] = dgst[3];
+    tmps[gid].out[i + 4] = dgst[4];
+    tmps[gid].out[i + 5] = dgst[5];
+    tmps[gid].out[i + 6] = dgst[6];
+    tmps[gid].out[i + 7] = dgst[7];
+  }
+}
+
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m13000_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha256_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global pbkdf2_sha256_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  const u32 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  u32 ipad[8];
+
+  ipad[0] = tmps[gid].ipad[0];
+  ipad[1] = tmps[gid].ipad[1];
+  ipad[2] = tmps[gid].ipad[2];
+  ipad[3] = tmps[gid].ipad[3];
+  ipad[4] = tmps[gid].ipad[4];
+  ipad[5] = tmps[gid].ipad[5];
+  ipad[6] = tmps[gid].ipad[6];
+  ipad[7] = tmps[gid].ipad[7];
+
+  u32 opad[8];
+
+  opad[0] = tmps[gid].opad[0];
+  opad[1] = tmps[gid].opad[1];
+  opad[2] = tmps[gid].opad[2];
+  opad[3] = tmps[gid].opad[3];
+  opad[4] = tmps[gid].opad[4];
+  opad[5] = tmps[gid].opad[5];
+  opad[6] = tmps[gid].opad[6];
+  opad[7] = tmps[gid].opad[7];
+
+  for (u32 i = 0; i < 8; i += 8)
+  {
+    u32 dgst[8];
+
+    dgst[0] = tmps[gid].dgst[i + 0];
+    dgst[1] = tmps[gid].dgst[i + 1];
+    dgst[2] = tmps[gid].dgst[i + 2];
+    dgst[3] = tmps[gid].dgst[i + 3];
+    dgst[4] = tmps[gid].dgst[i + 4];
+    dgst[5] = tmps[gid].dgst[i + 5];
+    dgst[6] = tmps[gid].dgst[i + 6];
+    dgst[7] = tmps[gid].dgst[i + 7];
+
+    u32 out[8];
+
+    out[0] = tmps[gid].out[i + 0];
+    out[1] = tmps[gid].out[i + 1];
+    out[2] = tmps[gid].out[i + 2];
+    out[3] = tmps[gid].out[i + 3];
+    out[4] = tmps[gid].out[i + 4];
+    out[5] = tmps[gid].out[i + 5];
+    out[6] = tmps[gid].out[i + 6];
+    out[7] = tmps[gid].out[i + 7];
+
+    for (u32 j = 0; j < loop_cnt; j++)
+    {
+      u32 w0[4];
+      u32 w1[4];
+      u32 w2[4];
+      u32 w3[4];
+
+      w0[0] = dgst[0];
+      w0[1] = dgst[1];
+      w0[2] = dgst[2];
+      w0[3] = dgst[3];
+      w1[0] = dgst[4];
+      w1[1] = dgst[5];
+      w1[2] = dgst[6];
+      w1[3] = dgst[7];
+      w2[0] = 0x80000000;
+      w2[1] = 0;
+      w2[2] = 0;
+      w2[3] = 0;
+      w3[0] = 0;
+      w3[1] = 0;
+      w3[2] = 0;
+      w3[3] = (64 + 32) * 8;
+
+      hmac_sha256_run (w0, w1, w2, w3, ipad, opad, dgst);
+
+      out[0] ^= dgst[0];
+      out[1] ^= dgst[1];
+      out[2] ^= dgst[2];
+      out[3] ^= dgst[3];
+      out[4] ^= dgst[4];
+      out[5] ^= dgst[5];
+      out[6] ^= dgst[6];
+      out[7] ^= dgst[7];
+    }
+
+    tmps[gid].dgst[i + 0] = dgst[0];
+    tmps[gid].dgst[i + 1] = dgst[1];
+    tmps[gid].dgst[i + 2] = dgst[2];
+    tmps[gid].dgst[i + 3] = dgst[3];
+    tmps[gid].dgst[i + 4] = dgst[4];
+    tmps[gid].dgst[i + 5] = dgst[5];
+    tmps[gid].dgst[i + 6] = dgst[6];
+    tmps[gid].dgst[i + 7] = dgst[7];
+
+    tmps[gid].out[i + 0] = out[0];
+    tmps[gid].out[i + 1] = out[1];
+    tmps[gid].out[i + 2] = out[2];
+    tmps[gid].out[i + 3] = out[3];
+    tmps[gid].out[i + 4] = out[4];
+    tmps[gid].out[i + 5] = out[5];
+    tmps[gid].out[i + 6] = out[6];
+    tmps[gid].out[i + 7] = out[7];
+  }
+}
+
+__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m13000_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha256_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global pbkdf2_sha256_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
+{
+  /**
+   * base
+   */
+
+  const u32 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  const u32 lid = get_local_id (0);
+
+  u32 outl4[4];
+
+  outl4[0] = tmps[gid].out[0];
+  outl4[1] = tmps[gid].out[2];
+  outl4[2] = tmps[gid].out[4];
+  outl4[3] = tmps[gid].out[6];
+
+  u32 outr4[4];
+
+  outr4[0] = tmps[gid].out[1];
+  outr4[1] = tmps[gid].out[3];
+  outr4[2] = tmps[gid].out[5];
+  outr4[3] = tmps[gid].out[7];
+
+  u32 outl = outl4[0]
+           ^ outl4[1]
+           ^ outl4[2]
+           ^ outl4[3];
+
+  u32 outr = outr4[0]
+           ^ outr4[1]
+           ^ outr4[2]
+           ^ outr4[3];
+
+  const u32 r0 = outl;
+  const u32 r1 = outr;
+  const u32 r2 = 0;
+  const u32 r3 = 0;
+
+  #define il_pos 0
+
+  #include COMPARE_M
+}

diff --git a/OpenCL/types_ocl.c b/OpenCL/types_ocl.c
index d71744d..5367a09 100644 (file)
--- a/OpenCL/types_ocl.c
+++ b/OpenCL/types_ocl.c
@@ -490,6 +490,8 @@ typedef struct
   u32  digest_buf[8];
   #elif defined _ANDROIDFDE_SAMSUNG_
   u32  digest_buf[8];
+  #elif defined _RAR5_
+  u32  digest_buf[4];
   #endif
 
 } digest_t;

diff --git a/docs/changes.txt b/docs/changes.txt
index 8a4b2b5..a46ea79 100644 (file)
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -13,6 +13,11 @@ File.: Kernel
 Desc.: Added new hash mode -m 12900 = Android FDE (Samsung DEK)
 Issue: 86
 
+Type.: Feature
+File.: Kernel
+Desc.: Added new hash mode -m 13000 = RAR5
+Issue: 68
+
 Type.: Feature
 File.: Host
 Desc.: Added a Makefile target for native compilation

diff --git a/docs/readme.txt b/docs/readme.txt
index 517facf..8e6d7ae 100644 (file)
--- a/docs/readme.txt
+++ b/docs/readme.txt
@@ -183,6 +183,7 @@ AMD users require Catalyst 14.9 or later (recommended 15.9 or later)
 - Skype
 - 7-Zip
 - RAR3-hp
+- RAR5
 - PDF 1.1 - 1.3 (Acrobat 2 - 4)
 - PDF 1.4 - 1.6 (Acrobat 5 - 8)
 - PDF 1.7 Level 3 (Acrobat 9)

diff --git a/extra/tab_completion/oclHashcat.sh b/extra/tab_completion/oclHashcat.sh
index 7f6a5d6..4910023 100644 (file)
--- a/extra/tab_completion/oclHashcat.sh
+++ b/extra/tab_completion/oclHashcat.sh
@@ -171,7 +171,7 @@ _oclHashcat ()
 {
   local VERSION=2.10
 
-  local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 190 200 300 400 500 501 900 1000 1100 1400 1410 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 4300 4400 4500 4700 4800 4900 5000 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12100 12200 12300 12400 12500 12600 12700 12800 12900"
+  local HASH_MODES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 124 130 131 132 133 140 141 150 160 190 200 300 400 500 501 900 1000 1100 1400 1410 1420 1421 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 4300 4400 4500 4700 4800 4900 5000 5100 5200 5300 5400 5500 5600 5700 5800 6000 6100 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8800 8900 9000 9100 9200 9300 9400 9500 9600 9700 9710 9720 9800 9810 9820 9900 10000 10100 10200 10300 10400 10410 10420 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11700 11800 11900 12000 12100 12200 12300 12400 12500 12600 12700 12800 12900 13000"
   local ATTACK_MODES="0 1 3 6 7"
   local OUTFILE_FORMATS="1 2 3 4 5 6 7 8 9 10 11 12 13 14 15"
   local BENCHMARK_MODE="0 1"

diff --git a/include/constants.h b/include/constants.h
index 74edf97..72671bf 100644 (file)
--- a/include/constants.h
+++ b/include/constants.h
@@ -71,7 +71,7 @@
 #define SHA1C03 0xca62c1d6u
 #endif
 
-#if defined _SHA256_ || defined _PDF17L8_ || defined _SEVEN_ZIP_ || defined _ANDROIDFDE_ || defined _CLOUDKEY_ || defined _SCRYPT_ || defined _PBKDF2_SHA256_ || defined _SHA256_SHA1_ || defined _MS_DRSR_ || defined _ANDROIDFDE_SAMSUNG_
+#if defined _SHA256_ || defined _PDF17L8_ || defined _SEVEN_ZIP_ || defined _ANDROIDFDE_ || defined _CLOUDKEY_ || defined _SCRYPT_ || defined _PBKDF2_SHA256_ || defined _SHA256_SHA1_ || defined _MS_DRSR_ || defined _ANDROIDFDE_SAMSUNG_ || defined _RAR5_
 /**
  * SHA256 Constants
  */

diff --git a/include/kernel_functions.c b/include/kernel_functions.c
index af9b0e1..7ef6fa5 100644 (file)
--- a/include/kernel_functions.c
+++ b/include/kernel_functions.c
@@ -193,7 +193,7 @@
 }
 #endif
 
-#if defined _SHA256_ || defined _PDF17L8_ || defined _SEVEN_ZIP_ || defined _ANDROIDFDE_ || defined _CLOUDKEY_ || defined _SCRYPT_ || defined _PBKDF2_SHA256_  || defined _SHA256_SHA1_ || defined _MS_DRSR_ || defined _ANDROIDFDE_SAMSUNG_
+#if defined _SHA256_ || defined _PDF17L8_ || defined _SEVEN_ZIP_ || defined _ANDROIDFDE_ || defined _CLOUDKEY_ || defined _SCRYPT_ || defined _PBKDF2_SHA256_  || defined _SHA256_SHA1_ || defined _MS_DRSR_ || defined _ANDROIDFDE_SAMSUNG_ || defined _RAR5_
 
 #define SHIFT_RIGHT_32(x,n) ((x) >> (n))
 

diff --git a/include/shared.h b/include/shared.h
index 1d4f1ce..919dcc7 100644 (file)
--- a/include/shared.h
+++ b/include/shared.h
@@ -352,6 +352,7 @@ extern hc_thread_mutex_t mux_display;
 #define KERNEL_ACCEL_12700   64
 #define KERNEL_ACCEL_12800   64
 #define KERNEL_ACCEL_12900   8
+#define KERNEL_ACCEL_13000   8
 
 #define KERNEL_LOOPS_0       256
 #define KERNEL_LOOPS_10      256
@@ -530,6 +531,7 @@ extern hc_thread_mutex_t mux_display;
 #define KERNEL_LOOPS_12700   10
 #define KERNEL_LOOPS_12800   100
 #define KERNEL_LOOPS_12900   64
+#define KERNEL_LOOPS_13000   64
 
 /**
  * Strings
@@ -676,6 +678,7 @@ extern hc_thread_mutex_t mux_display;
 #define HT_12700  "Blockchain, My Wallet"
 #define HT_12800  "MS-AzureSync PBKDF2-HMAC-SHA256"
 #define HT_12900  "Android FDE (Samsung DEK)"
+#define HT_13000  "RAR5"
 
 #define HT_00011  "Joomla < 2.5.18"
 #define HT_00012  "PostgreSQL"
@@ -1007,6 +1010,8 @@ extern hc_thread_mutex_t mux_display;
 #define DISPLAY_LEN_MAX_12800 11 + 1 + 20 + 1 + 5 + 1 + 64
 #define DISPLAY_LEN_MIN_12900 64 + 64 + 32
 #define DISPLAY_LEN_MAX_12900 64 + 64 + 32
+#define DISPLAY_LEN_MIN_13000 1 + 4 + 1 + 2 + 1 + 32 + 1 + 2 + 1 + 32 + 1 + 1 + 1 + 16
+#define DISPLAY_LEN_MAX_13000 1 + 4 + 1 + 2 + 1 + 32 + 1 + 2 + 1 + 32 + 1 + 1 + 1 + 16
 
 #define DISPLAY_LEN_MIN_11    32 + 1 + 16
 #define DISPLAY_LEN_MAX_11    32 + 1 + 32
@@ -1264,6 +1269,7 @@ extern hc_thread_mutex_t mux_display;
 #define KERN_TYPE_MYWALLET        12700
 #define KERN_TYPE_MS_DRSR         12800
 #define KERN_TYPE_ANDROIDFDE_SAMSUNG  12900
+#define KERN_TYPE_RAR5            13000
 
 /**
  * signatures
@@ -1332,6 +1338,7 @@ extern hc_thread_mutex_t mux_display;
 #define SIGNATURE_RAR3            "$RAR3$"
 #define SIGNATURE_MYWALLET        "$blockchain$"
 #define SIGNATURE_MS_DRSR         "v1;PPH1_MD4"
+#define SIGNATURE_RAR5            "$rar5$"
 
 /**
  * Default iteration numbers
@@ -1382,6 +1389,7 @@ extern hc_thread_mutex_t mux_display;
 #define ROUNDS_MYWALLET       10
 #define ROUNDS_MS_DRSR        100
 #define ROUNDS_ANDROIDFDE_SAMSUNG 4096
+#define ROUNDS_RAR5           (1 << 15)
 
 /**
  * salt types
@@ -1861,6 +1869,7 @@ int pbkdf2_sha512_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
 int ecryptfs_parse_hash     (char *input_buf, uint input_len, hash_t *hash_buf);
 int bsdicrypt_parse_hash    (char *input_buf, uint input_len, hash_t *hash_buf);
 int rar3hp_parse_hash       (char *input_buf, uint input_len, hash_t *hash_buf);
+int rar5_parse_hash         (char *input_buf, uint input_len, hash_t *hash_buf);
 int cf10_parse_hash         (char *input_buf, uint input_len, hash_t *hash_buf);
 int mywallet_parse_hash     (char *input_buf, uint input_len, hash_t *hash_buf);
 int ms_drsr_parse_hash      (char *input_buf, uint input_len, hash_t *hash_buf);

diff --git a/src/oclHashcat.c b/src/oclHashcat.c
index 510266a..82a5211 100644 (file)
--- a/src/oclHashcat.c
+++ b/src/oclHashcat.c
@@ -128,7 +128,7 @@ const uint  RESTORE_MIN       = 210;
 
 #define MAX_DICTSTAT            10000
 
-#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 129
+#define NUM_DEFAULT_BENCHMARK_ALGORITHMS 130
 
 #define global_free(attr)       \
 {                               \
@@ -249,6 +249,7 @@ static uint default_benchmark_algorithms[NUM_DEFAULT_BENCHMARK_ALGORITHMS] =
   133,
   11600,
   12500,
+  13000,
   6211,
   6221,
   6231,
@@ -668,6 +669,7 @@ const char *USAGE_BIG[] =
   "",
   "  11600 = 7-Zip",
   "  12500 = RAR3-hp",
+  "  13000 = RAR5",
   "",
   "[[ Full-Disk encryptions (FDE) ]]",
   "",
@@ -5610,7 +5612,7 @@ int main (int argc, char **argv)
     return (-1);
   }
 
-  if (hash_mode_chgd && hash_mode > 12900) // just added to remove compiler warnings for hash_mode_chgd
+  if (hash_mode_chgd && hash_mode > 13000) // just added to remove compiler warnings for hash_mode_chgd
   {
     log_error ("ERROR: Invalid hash-type specified");
 
@@ -9792,6 +9794,21 @@ int main (int argc, char **argv)
                    dgst_pos3   = 3;
                    break;
 
+      case 13000:  hash_type   = HASH_TYPE_PBKDF2_SHA256;
+                   salt_type   = SALT_TYPE_EMBEDDED;
+                   attack_exec = ATTACK_EXEC_OUTSIDE_KERNEL;
+                   opts_type   = OPTS_TYPE_PT_GENERATE_LE;
+                   kern_type   = KERN_TYPE_RAR5;
+                   dgst_size   = DGST_SIZE_4_4;
+                   parse_func  = rar5_parse_hash;
+                   sort_by_digest = sort_by_digest_4_4;
+                   opti_type   = OPTI_TYPE_ZERO_BYTE;
+                   dgst_pos0   = 0;
+                   dgst_pos1   = 1;
+                   dgst_pos2   = 2;
+                   dgst_pos3   = 3;
+                   break;
+
       default:     usage_mini_print (PROGNAME); return (-1);
     }
 
@@ -11148,6 +11165,8 @@ int main (int argc, char **argv)
                      break;
         case 12900:  hashes_buf[0].salt->salt_iter = ROUNDS_ANDROIDFDE_SAMSUNG - 1;
                      break;
+        case 13000:  hashes_buf[0].salt->salt_iter = ROUNDS_RAR5 - 1;
+                     break;
       }
 
       // set special tuning for benchmark-mode 1
@@ -11333,6 +11352,9 @@ int main (int argc, char **argv)
           case 12900:  kernel_loops = ROUNDS_ANDROIDFDE_SAMSUNG;
                        kernel_accel = 8;
                        break;
+          case 13000:  kernel_loops = ROUNDS_RAR5;
+                       kernel_accel = 8;
+                       break;
         }
 
         // some algorithm collide too fast, make that impossible
@@ -13112,6 +13134,7 @@ int main (int argc, char **argv)
         case 12700: size_tmps = kernel_blocks * sizeof (mywallet_tmp_t);      break;
         case 12800: size_tmps = kernel_blocks * sizeof (pbkdf2_sha256_tmp_t); break;
         case 12900: size_tmps = kernel_blocks * sizeof (pbkdf2_sha256_tmp_t); break;
+        case 13000: size_tmps = kernel_blocks * sizeof (pbkdf2_sha256_tmp_t); break;
       };
 
       uint size_hooks = 4;

diff --git a/src/shared.c b/src/shared.c
index 7791b93..9b251f1 100644 (file)
--- a/src/shared.c
+++ b/src/shared.c
@@ -5611,6 +5611,7 @@ char *strhashtype (const uint hash_mode)
     case 12700: return ((char *) HT_12700); break;
     case 12800: return ((char *) HT_12800); break;
     case 12900: return ((char *) HT_12900); break;
+    case 13000: return ((char *) HT_13000); break;
   }
 
   return ((char *) "Unknown");
@@ -8068,6 +8069,22 @@ void ascii_digest (char out_buf[4096], uint salt_pos, uint digest_pos)
       salt.salt_buf[ 3]
     );
   }
+  else if (hash_mode == 13000)
+  {
+    snprintf (out_buf, len-1, "$rar5$16$%08x%08x%08x%08x$%u$%08x%08x%08x%08x$8$%08x%08x",
+      salt.salt_buf[0],
+      salt.salt_buf[1],
+      salt.salt_buf[2],
+      salt.salt_buf[3],
+      salt.salt_sign[0],
+      salt.salt_buf[4],
+      salt.salt_buf[5],
+      salt.salt_buf[6],
+      salt.salt_buf[7],
+      byte_swap_32 (digest_buf[0]),
+      byte_swap_32 (digest_buf[1])
+    );
+  }
   else
   {
     if (hash_type == HASH_TYPE_MD4)
@@ -9023,6 +9040,7 @@ uint set_kernel_accel (uint hash_mode)
     case 12700: return GET_ACCEL (12700);
     case 12800: return GET_ACCEL (12800);
     case 12900: return GET_ACCEL (12900);
+    case 13000: return GET_ACCEL (13000);
   }
 
   return 0;
@@ -9209,6 +9227,7 @@ uint set_kernel_loops (uint hash_mode)
     case 12700: return GET_LOOPS (12700);
     case 12800: return GET_LOOPS (12800);
     case 12900: return GET_LOOPS (12900);
+    case 13000: return GET_LOOPS (13000);
   }
 
   return 0;
@@ -18365,6 +18384,116 @@ int rar3hp_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
   return (PARSER_OK);
 }
 
+int rar5_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
+{
+  if ((input_len < DISPLAY_LEN_MIN_13000) || (input_len > DISPLAY_LEN_MAX_13000)) return (PARSER_GLOBAL_LENGTH);
+
+  if (memcmp (SIGNATURE_RAR5, input_buf, 1 + 4 + 1)) return (PARSER_SIGNATURE_UNMATCHED);
+
+  uint32_t *digest = (uint32_t *) hash_buf->digest;
+
+  salt_t *salt = hash_buf->salt;
+
+  /**
+   * parse line
+   */
+
+  char *param0_pos = input_buf + 1 + 4 + 1;
+
+  char *param1_pos = strchr (param0_pos, '$');
+
+  if (param1_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  uint32_t param0_len = param1_pos - param0_pos;
+
+  param1_pos++;
+
+  char *param2_pos = strchr (param1_pos, '$');
+
+  if (param2_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  uint32_t param1_len = param2_pos - param1_pos;
+
+  param2_pos++;
+
+  char *param3_pos = strchr (param2_pos, '$');
+
+  if (param3_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  uint32_t param2_len = param3_pos - param2_pos;
+
+  param3_pos++;
+
+  char *param4_pos = strchr (param3_pos, '$');
+
+  if (param4_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  uint32_t param3_len = param4_pos - param3_pos;
+
+  param4_pos++;
+
+  char *param5_pos = strchr (param4_pos, '$');
+
+  if (param5_pos == NULL) return (PARSER_SEPARATOR_UNMATCHED);
+
+  uint32_t param4_len = param5_pos - param4_pos;
+
+  param5_pos++;
+
+  uint32_t param5_len = input_len - 1 - 4 - 1 - param0_len - 1 - param1_len - 1 - param2_len - 1 - param3_len - 1 - param4_len - 1;
+
+  char *salt_buf = param1_pos;
+  char *iv       = param3_pos;
+  char *pswcheck = param5_pos;
+
+  const uint salt_len     = atoi (param0_pos);
+  const uint iterations   = atoi (param2_pos);
+  const uint pswcheck_len = atoi (param4_pos);
+
+  /**
+   * verify some data
+   */
+
+  if (param1_len   != 32) return (PARSER_SALT_VALUE);
+  if (param3_len   != 32) return (PARSER_SALT_VALUE);
+  if (param5_len   != 16) return (PARSER_SALT_VALUE);
+
+  if (salt_len     != 16) return (PARSER_SALT_VALUE);
+  if (iterations   ==  0) return (PARSER_SALT_VALUE);
+  if (pswcheck_len !=  8) return (PARSER_SALT_VALUE);
+
+  /**
+   * store data
+   */
+
+  salt->salt_buf[0] = hex_to_uint (&salt_buf[ 0]);
+  salt->salt_buf[1] = hex_to_uint (&salt_buf[ 8]);
+  salt->salt_buf[2] = hex_to_uint (&salt_buf[16]);
+  salt->salt_buf[3] = hex_to_uint (&salt_buf[24]);
+
+  salt->salt_buf[4] = hex_to_uint (&iv[ 0]);
+  salt->salt_buf[5] = hex_to_uint (&iv[ 8]);
+  salt->salt_buf[6] = hex_to_uint (&iv[16]);
+  salt->salt_buf[7] = hex_to_uint (&iv[24]);
+
+  salt->salt_len = 16 + 16;
+
+  salt->salt_sign[0] = iterations;
+
+  salt->salt_iter = ((1 << iterations) + 32) - 1;
+
+  /**
+   * digest buf
+   */
+
+  digest[0] = hex_to_uint (&pswcheck[ 0]);
+  digest[1] = hex_to_uint (&pswcheck[ 8]);
+  digest[2] = 0;
+  digest[3] = 0;
+
+  return (PARSER_OK);
+}
+
 int cf10_parse_hash (char *input_buf, uint input_len, hash_t *hash_buf)
 {
   if ((input_len < DISPLAY_LEN_MIN_12600) || (input_len > DISPLAY_LEN_MAX_12600)) return (PARSER_GLOBAL_LENGTH);

diff --git a/tools/test.pl b/tools/test.pl
index 654175e..f02d707 100755 (executable)
--- a/tools/test.pl
+++ b/tools/test.pl
@@ -44,7 +44,7 @@ my $hashcat = "./oclHashcat";
 
 my $MAX_LEN = 55;
 
-my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 130, 131, 132, 140, 141, 150, 160, 190, 200, 300, 400, 500, 900, 1000, 1100, 1400, 1410, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7100, 7200, 7300, 7400, 7500, 7600, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900);
+my @modes = (0, 10, 11, 12, 20, 21, 22, 23, 30, 40, 50, 60, 100, 101, 110, 111, 112, 120, 121, 122, 130, 131, 132, 140, 141, 150, 160, 190, 200, 300, 400, 500, 900, 1000, 1100, 1400, 1410, 1420, 1430, 1440, 1441, 1450, 1460, 1500, 1600, 1700, 1710, 1711, 1720, 1730, 1740, 1722, 1731, 1750, 1760, 1800, 2100, 2400, 2410, 2500, 2600, 2611, 2612, 2711, 2811, 3000, 3100, 3200, 3710, 3711, 3300, 3500, 3610, 3720, 3800, 3910, 4010, 4110, 4210, 4300, 4400, 4500, 4600, 4700, 4800, 4900, 5000, 5100, 5300, 5400, 5500, 5600, 5700, 5800, 6000, 6100, 6300, 6400, 6500, 6600, 6700, 6800, 6900, 7100, 7200, 7300, 7400, 7500, 7600, 7700, 7800, 7900, 8000, 8100, 8200, 8300, 8400, 8500, 8600, 8700, 8900, 9100, 9200, 9300, 9400, 9500, 9600, 9700, 9800, 9900, 10000, 10100, 10200, 10300, 10400, 10500, 10600, 10700, 10800, 10900, 11000, 11100, 11200, 11300, 11400, 11500, 11600, 11900, 12000, 12100, 12200, 12300, 12400, 12600, 12700, 12800, 12900, 13000);
 
 my %is_unicode      = map { $_ => 1 } qw(30 40 130 131 132 140 141 1000 1100 1430 1440 1441 1730 1740 1731 5500 5600 8000 9400 9500 9600 9700 9800);
 my %less_fifteen    = map { $_ => 1 } qw(500 1600 1800 2400 2410 3200 6300 7400 10500 10700);
@@ -2175,6 +2175,40 @@ sub verify
 
       next unless (exists ($db->{$hash_in}) and (! defined ($db->{$hash_in})));
     }
+    elsif ($mode == 13000)
+    {
+      my $hash_line;
+
+      ($hash_line, $word) = split ":", $line;
+
+      next unless defined $hash_line;
+      next unless defined $word;
+
+      my @data = split ('\$', $hash_line);
+
+      next unless scalar @data == 8;
+
+      shift @data;
+
+      my $signature    = shift @data;
+      my $salt_len     = shift @data;
+      my $salt_buf     = shift @data;
+      my $iterations   = shift @data;
+      my $iv           = shift @data;
+      my $pswcheck_len = shift @data;
+      my $pswcheck     = shift @data;
+
+      next unless ($signature eq "rar5");
+      next unless ($salt_len == 16);
+      next unless ($pswcheck_len == 8);
+
+      $salt    = $salt_buf;
+      $iter    = $iterations;
+      $hash_in = $pswcheck;
+      $param   = $iv;
+
+      next unless (exists ($db->{$hash_line}) and (! defined ($db->{$hash_line})));
+    }
     else
     {
       print "ERROR: hash mode is not supported\n";
@@ -2449,6 +2483,14 @@ sub verify
 
       return unless (substr ($line, 0, $len) eq $hash_out);
     }
+    elsif ($mode == 13000)
+    {
+      $hash_out = gen_hash ($mode, $word, $salt, $iter, $param);
+
+      $len = length $hash_out;
+
+      return unless (substr ($line, 0, $len) eq $hash_out);
+    }
     else
     {
       $hash_out = gen_hash ($mode, $word, $salt, $iter);
@@ -2865,6 +2907,10 @@ sub passthrough
     {
       $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32));
     }
+    elsif ($mode == 13000)
+    {
+      $tmp_hash = gen_hash ($mode, $word_buf, substr ($salt_buf, 0, 32));
+    }
     else
     {
       print "ERROR: Unsupported hash type\n";
@@ -3578,6 +3624,20 @@ sub single
         }
       }
     }
+    elsif ($mode == 13000)
+    {
+      for (my $i = 1; $i < 32; $i++)
+      {
+        if ($len != 0)
+        {
+          rnd ($mode, $len, 32);
+        }
+        else
+        {
+          rnd ($mode, $i, 32);
+        }
+      }
+    }
   }
 }
 
@@ -6518,6 +6578,42 @@ END_CODE
 
     $tmp_hash = sprintf ("%s%s%s", $salt2, $hash_hmac, $salt_buf);
   }
+  elsif ($mode == 13000)
+  {
+    my $iterations = 15;
+
+    if (length ($iter))
+    {
+      $iterations = int ($iter);
+    }
+
+    my $iv = "0" x 32;
+
+    if (defined $additional_param)
+    {
+      $iv = $additional_param;
+    }
+
+    my $pbkdf2 = Crypt::PBKDF2->new
+    (
+      hasher     => Crypt::PBKDF2->hasher_from_algorithm ('HMACSHA2', 256),
+      iterations => (1 << $iterations) + 32,
+      output_len => 32
+    );
+
+    my $salt_buf_bin = pack ("H*", $salt_buf);
+
+    my $hash = $pbkdf2->PBKDF2 ($salt_buf_bin, $word_buf);
+
+printf ("%s\n", unpack ("H*", $hash));
+
+    my $hash_final = substr ($hash,  0, 8)
+                   ^ substr ($hash,  8, 8)
+                   ^ substr ($hash, 16, 8)
+                   ^ substr ($hash, 24, 8);
+
+    $tmp_hash = sprintf ('$rar5$16$%s$%d$%s$8$%s', $salt_buf, $iterations, $iv, unpack ("H*", $hash_final));
+  }
 
   return ($tmp_hash);
 }

diff --git a/tools/test.sh b/tools/test.sh
index f5259fb..7406796 100755 (executable)
--- a/tools/test.sh
+++ b/tools/test.sh
@@ -7,7 +7,7 @@
 
 # missing hash types: 5200,6211,6221,6231,6241,6251,6261,6271,6281
 
-HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 130 131 132 140 141 150 160 190 200 300 400 500 900 1000 1100 1400 1410 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 4300 4400 4500 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12100 12200 12300 12400 12600 12800 12900"
+HASH_TYPES="0 10 11 12 20 21 22 23 30 40 50 60 100 101 110 111 112 120 121 122 130 131 132 140 141 150 160 190 200 300 400 500 900 1000 1100 1400 1410 1420 1430 1440 1441 1450 1460 1500 1600 1700 1710 1711 1720 1722 1730 1731 1740 1750 1760 1800 2100 2400 2410 2500 2600 2611 2612 2711 2811 3000 3100 3200 3710 3711 3800 4300 4400 4500 4700 4800 4900 5000 5100 5300 5400 5500 5600 5700 5800 6000 6100 6300 6400 6500 6600 6700 6800 6900 7100 7200 7300 7400 7500 7600 7700 7800 7900 8000 8100 8200 8300 8400 8500 8600 8700 8900 9100 9200 9300 9400 9500 9600 9700 9800 9900 10000 10100 10200 10300 10400 10500 10600 10700 10800 10900 11000 11100 11200 11300 11400 11500 11600 11900 12000 12100 12200 12300 12400 12600 12800 12900 13000"
 
 ATTACK_MODES="0 1 3 6 7"
 
@@ -17,7 +17,7 @@ HASHFILE_ONLY="2500"
 
 NEVER_CRACK="11600"
 
-SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6221 6231 6241 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12100 12200 12300 12400 12500 12800 12900"
+SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6221 6231 6241 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12100 12200 12300 12400 12500 12800 12900 13000"
 
 OPTS="--quiet --force --potfile-disable --runtime 200 --gpu-temp-disable -d 1 --weak-hash-threshold=0"