Added hash-mode 18700 DJB 32

OpenCL/m18700_a0-optimized.cl

@@ -0,0 +1,165 @@
+/**
+ * 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_optimized.h"
+#include "inc_rp_optimized.cl"
+#include "inc_simd.cl"
+
+__kernel void m18700_m04 (KERN_ATTR_RULES ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+
+  /**
+   * base
+   */
+
+  const u64 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 & 63;
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    u32x w[16] = { 0 };
+
+    const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w + 0, w + 4);
+
+    u32x hash = 0;
+
+    for (u32 i = 0; i < pw_len; i++)
+    {
+      const u32 c32 = w[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_M_SIMD (r0, r1, r2, r3);
+  }
+}
+
+__kernel void m18700_m08 (KERN_ATTR_RULES ())
+{
+}
+
+__kernel void m18700_m16 (KERN_ATTR_RULES ())
+{
+}
+
+__kernel void m18700_s04 (KERN_ATTR_RULES ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+
+  /**
+   * base
+   */
+
+  const u64 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 & 63;
+
+  /**
+   * digest
+   */
+
+  const u32 search[4] =
+  {
+    digests_buf[digests_offset].digest_buf[DGST_R0],
+    0,
+    0,
+    0
+  };
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    u32x w[16] = { 0 };
+
+    const u32x out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w + 0, w + 4);
+
+    u32x hash = 0;
+
+    for (u32 i = 0; i < pw_len; i++)
+    {
+      const u32 c32 = w[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_S_SIMD (r0, r1, r2, r3);
+  }
+}
+
+__kernel void m18700_s08 (KERN_ATTR_RULES ())
+{
+}
+
+__kernel void m18700_s16 (KERN_ATTR_RULES ())
+{
+}

OpenCL/m18700_a0-pure.cl

@@ -0,0 +1,124 @@
+/**
+ * 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_scalar.cl"
+#include "inc_hash_md5.cl"
+
+__kernel void m18700_mxx (KERN_ATTR_RULES ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  COPY_PW (pws[gid]);
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
+  {
+    pw_t tmp = PASTE_PW;
+
+    tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
+
+    u32 hash = 0;
+
+    for (u32 i = 0; i < tmp.pw_len; i++)
+    {
+      const u32 c32 = tmp.i[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32 r0 = hash;
+    const u32 r1 = 0;
+    const u32 r2 = 0;
+    const u32 r3 = 0;
+
+    COMPARE_M_SCALAR (r0, r1, r2, r3);
+  }
+}
+
+__kernel void m18700_sxx (KERN_ATTR_RULES ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * digest
+   */
+
+  const u32 search[4] =
+  {
+    digests_buf[digests_offset].digest_buf[DGST_R0],
+    0,
+    0,
+    0
+  };
+
+  /**
+   * base
+   */
+
+  COPY_PW (pws[gid]);
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
+  {
+    pw_t tmp = PASTE_PW;
+
+    tmp.pw_len = apply_rules (rules_buf[il_pos].cmds, tmp.i, tmp.pw_len);
+
+    u32 hash = 0;
+
+    for (u32 i = 0; i < tmp.pw_len; i++)
+    {
+      const u32 c32 = tmp.i[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32 r0 = hash;
+    const u32 r1 = 0;
+    const u32 r2 = 0;
+    const u32 r3 = 0;
+
+    COMPARE_S_SCALAR (r0, r1, r2, r3);
+  }
+}

OpenCL/m18700_a1-optimized.cl

@@ -0,0 +1,293 @@
+/**
+ * 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_scalar.cl"
+#include "inc_simd.cl"
+#include "inc_hash_md5.cl"
+
+__kernel void m18700_m04 (KERN_ATTR_BASIC ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+
+  /**
+   * base
+   */
+
+  const u64 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_l_len = pws[gid].pw_len & 63;
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63;
+
+    const u32x pw_len = (pw_l_len + pw_r_len) & 63;
+
+    /**
+     * concat password candidate
+     */
+
+    u32x wordl0[4] = { 0 };
+    u32x wordl1[4] = { 0 };
+    u32x wordl2[4] = { 0 };
+    u32x wordl3[4] = { 0 };
+
+    wordl0[0] = pw_buf0[0];
+    wordl0[1] = pw_buf0[1];
+    wordl0[2] = pw_buf0[2];
+    wordl0[3] = pw_buf0[3];
+    wordl1[0] = pw_buf1[0];
+    wordl1[1] = pw_buf1[1];
+    wordl1[2] = pw_buf1[2];
+    wordl1[3] = pw_buf1[3];
+
+    u32x wordr0[4] = { 0 };
+    u32x wordr1[4] = { 0 };
+    u32x wordr2[4] = { 0 };
+    u32x wordr3[4] = { 0 };
+
+    wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
+    wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
+    wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
+    wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
+    wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
+    wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
+    wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
+    wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
+
+    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
+    {
+      switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
+    }
+    else
+    {
+      switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
+    }
+
+    u32x w[16];
+
+    w[ 0] = wordl0[0] | wordr0[0];
+    w[ 1] = wordl0[1] | wordr0[1];
+    w[ 2] = wordl0[2] | wordr0[2];
+    w[ 3] = wordl0[3] | wordr0[3];
+    w[ 4] = wordl1[0] | wordr1[0];
+    w[ 5] = wordl1[1] | wordr1[1];
+    w[ 6] = wordl1[2] | wordr1[2];
+    w[ 7] = wordl1[3] | wordr1[3];
+    w[ 8] = wordl2[0] | wordr2[0];
+    w[ 9] = wordl2[1] | wordr2[1];
+    w[10] = wordl2[2] | wordr2[2];
+    w[11] = wordl2[3] | wordr2[3];
+    w[12] = wordl3[0] | wordr3[0];
+    w[13] = wordl3[1] | wordr3[1];
+    w[14] = wordl3[2] | wordr3[2];
+    w[15] = wordl3[3] | wordr3[3];
+
+    /**
+     * djb
+     */
+
+    u32x hash = 0;
+
+    for (u32 i = 0; i < pw_len; i++)
+    {
+      const u32 c32 = w[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_M_SIMD (r0, r1, r2, r3);
+  }
+}
+
+__kernel void m18700_m08 (KERN_ATTR_BASIC ())
+{
+}
+
+__kernel void m18700_m16 (KERN_ATTR_BASIC ())
+{
+}
+
+__kernel void m18700_s04 (KERN_ATTR_BASIC ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+
+  /**
+   * base
+   */
+
+  const u64 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_l_len = pws[gid].pw_len & 63;
+
+  /**
+   * digest
+   */
+
+  const u32 search[4] =
+  {
+    digests_buf[digests_offset].digest_buf[DGST_R0],
+    0,
+    0,
+    0
+  };
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63;
+
+    const u32x pw_len = (pw_l_len + pw_r_len) & 63;
+
+    /**
+     * concat password candidate
+     */
+
+    u32x wordl0[4] = { 0 };
+    u32x wordl1[4] = { 0 };
+    u32x wordl2[4] = { 0 };
+    u32x wordl3[4] = { 0 };
+
+    wordl0[0] = pw_buf0[0];
+    wordl0[1] = pw_buf0[1];
+    wordl0[2] = pw_buf0[2];
+    wordl0[3] = pw_buf0[3];
+    wordl1[0] = pw_buf1[0];
+    wordl1[1] = pw_buf1[1];
+    wordl1[2] = pw_buf1[2];
+    wordl1[3] = pw_buf1[3];
+
+    u32x wordr0[4] = { 0 };
+    u32x wordr1[4] = { 0 };
+    u32x wordr2[4] = { 0 };
+    u32x wordr3[4] = { 0 };
+
+    wordr0[0] = ix_create_combt (combs_buf, il_pos, 0);
+    wordr0[1] = ix_create_combt (combs_buf, il_pos, 1);
+    wordr0[2] = ix_create_combt (combs_buf, il_pos, 2);
+    wordr0[3] = ix_create_combt (combs_buf, il_pos, 3);
+    wordr1[0] = ix_create_combt (combs_buf, il_pos, 4);
+    wordr1[1] = ix_create_combt (combs_buf, il_pos, 5);
+    wordr1[2] = ix_create_combt (combs_buf, il_pos, 6);
+    wordr1[3] = ix_create_combt (combs_buf, il_pos, 7);
+
+    if (combs_mode == COMBINATOR_MODE_BASE_LEFT)
+    {
+      switch_buffer_by_offset_le_VV (wordr0, wordr1, wordr2, wordr3, pw_l_len);
+    }
+    else
+    {
+      switch_buffer_by_offset_le_VV (wordl0, wordl1, wordl2, wordl3, pw_r_len);
+    }
+
+    u32x w[16];
+
+    w[ 0] = wordl0[0] | wordr0[0];
+    w[ 1] = wordl0[1] | wordr0[1];
+    w[ 2] = wordl0[2] | wordr0[2];
+    w[ 3] = wordl0[3] | wordr0[3];
+    w[ 4] = wordl1[0] | wordr1[0];
+    w[ 5] = wordl1[1] | wordr1[1];
+    w[ 6] = wordl1[2] | wordr1[2];
+    w[ 7] = wordl1[3] | wordr1[3];
+    w[ 8] = wordl2[0] | wordr2[0];
+    w[ 9] = wordl2[1] | wordr2[1];
+    w[10] = wordl2[2] | wordr2[2];
+    w[11] = wordl2[3] | wordr2[3];
+    w[12] = wordl3[0] | wordr3[0];
+    w[13] = wordl3[1] | wordr3[1];
+    w[14] = wordl3[2] | wordr3[2];
+    w[15] = wordl3[3] | wordr3[3];
+
+    /**
+     * djb
+     */
+
+    u32x hash = 0;
+
+    for (u32 i = 0; i < pw_len; i++)
+    {
+      const u32 c32 = w[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_S_SIMD (r0, r1, r2, r3);
+  }
+}
+
+__kernel void m18700_s08 (KERN_ATTR_BASIC ())
+{
+}
+
+__kernel void m18700_s16 (KERN_ATTR_BASIC ())
+{
+}

OpenCL/m18700_a1-pure.cl

@@ -0,0 +1,134 @@
+/**
+ * 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_scalar.cl"
+#include "inc_hash_md5.cl"
+
+__kernel void m18700_mxx (KERN_ATTR_BASIC ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  u32 base = 0;
+
+  for (u32 i = 0; i < pws[gid].pw_len; i++)
+  {
+    const u32 c32 = pws[gid].i[i / 4];
+
+    const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+    base *= 31;
+    base += c;
+  }
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
+  {
+    u32 hash = base;
+
+    for (u32 i = 0; i < combs_buf[il_pos].pw_len; i++)
+    {
+      const u32 c32 = combs_buf[il_pos].i[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_M_SCALAR (r0, r1, r2, r3);
+  }
+}
+
+__kernel void m18700_sxx (KERN_ATTR_BASIC ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * digest
+   */
+
+  const u32 search[4] =
+  {
+    digests_buf[digests_offset].digest_buf[DGST_R0],
+    0,
+    0,
+    0
+  };
+
+  /**
+   * base
+   */
+
+  u32 base = 0;
+
+  for (u32 i = 0; i < pws[gid].pw_len; i++)
+  {
+    const u32 c32 = pws[gid].i[i / 4];
+
+    const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+    base *= 31;
+    base += c;
+  }
+
+  /**
+   * loop
+   */
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos++)
+  {
+    u32 hash = base;
+
+    for (u32 i = 0; i < combs_buf[il_pos].pw_len; i++)
+    {
+      const u32 c32 = combs_buf[il_pos].i[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_S_SCALAR (r0, r1, r2, r3);
+  }
+}

OpenCL/m18700_a3-optimized.cl

@@ -0,0 +1,358 @@
+/**
+ * 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_simd.cl"
+#include "inc_hash_md5.cl"
+
+DECLSPEC void m18700m (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 gid = get_global_id (0);
+  const u64 lid = get_local_id (0);
+
+  /**
+   * loop
+   */
+
+  u32 w0l = w[0];
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
+
+    const u32x w0 = w0l | w0r;
+
+    u32x hash = 0;
+
+    for (u32 i = 0; i < 4; i++)
+    {
+      if (i == pw_len) break;
+
+      const u32x c = (w0 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    for (u32 i = 4; i < pw_len; i++)
+    {
+      const u32 c32 = w[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_M_SIMD (r0, r1, r2, r3);
+  }
+}
+
+DECLSPEC void m18700s (u32 *w, const u32 pw_len, KERN_ATTR_VECTOR ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 gid = get_global_id (0);
+  const u64 lid = get_local_id (0);
+
+  /**
+   * digest
+   */
+
+  const u32 search[4] =
+  {
+    digests_buf[digests_offset].digest_buf[DGST_R0],
+    0,
+    0,
+    0
+  };
+
+  /**
+   * loop
+   */
+
+  u32 w0l = w[0];
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
+
+    const u32x w0 = w0l | w0r;
+
+    u32x hash = 0;
+
+    for (u32 i = 0; i < 4; i++)
+    {
+      if (i == pw_len) break;
+
+      const u32x c = (w0 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    for (u32 i = 4; i < pw_len; i++)
+    {
+      const u32 c32 = w[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_S_SIMD (r0, r1, r2, r3);
+  }
+}
+
+__kernel void m18700_m04 (KERN_ATTR_VECTOR ())
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  u32 w[16];
+
+  w[ 0] = pws[gid].i[ 0];
+  w[ 1] = pws[gid].i[ 1];
+  w[ 2] = pws[gid].i[ 2];
+  w[ 3] = pws[gid].i[ 3];
+  w[ 4] = 0;
+  w[ 5] = 0;
+  w[ 6] = 0;
+  w[ 7] = 0;
+  w[ 8] = 0;
+  w[ 9] = 0;
+  w[10] = 0;
+  w[11] = 0;
+  w[12] = 0;
+  w[13] = 0;
+  w[14] = pws[gid].i[14];
+  w[15] = 0;
+
+  const u32 pw_len = pws[gid].pw_len & 63;
+
+  /**
+   * main
+   */
+
+  m18700m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
+}
+
+__kernel void m18700_m08 (KERN_ATTR_VECTOR ())
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  u32 w[16];
+
+  w[ 0] = pws[gid].i[ 0];
+  w[ 1] = pws[gid].i[ 1];
+  w[ 2] = pws[gid].i[ 2];
+  w[ 3] = pws[gid].i[ 3];
+  w[ 4] = pws[gid].i[ 4];
+  w[ 5] = pws[gid].i[ 5];
+  w[ 6] = pws[gid].i[ 6];
+  w[ 7] = pws[gid].i[ 7];
+  w[ 8] = 0;
+  w[ 9] = 0;
+  w[10] = 0;
+  w[11] = 0;
+  w[12] = 0;
+  w[13] = 0;
+  w[14] = pws[gid].i[14];
+  w[15] = 0;
+
+  const u32 pw_len = pws[gid].pw_len & 63;
+
+  /**
+   * main
+   */
+
+  m18700m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
+}
+
+__kernel void m18700_m16 (KERN_ATTR_VECTOR ())
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  u32 w[16];
+
+  w[ 0] = pws[gid].i[ 0];
+  w[ 1] = pws[gid].i[ 1];
+  w[ 2] = pws[gid].i[ 2];
+  w[ 3] = pws[gid].i[ 3];
+  w[ 4] = pws[gid].i[ 4];
+  w[ 5] = pws[gid].i[ 5];
+  w[ 6] = pws[gid].i[ 6];
+  w[ 7] = pws[gid].i[ 7];
+  w[ 8] = pws[gid].i[ 8];
+  w[ 9] = pws[gid].i[ 9];
+  w[10] = pws[gid].i[10];
+  w[11] = pws[gid].i[11];
+  w[12] = pws[gid].i[12];
+  w[13] = pws[gid].i[13];
+  w[14] = pws[gid].i[14];
+  w[15] = pws[gid].i[15];
+
+  const u32 pw_len = pws[gid].pw_len & 63;
+
+  /**
+   * main
+   */
+
+  m18700m (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
+}
+
+__kernel void m18700_s04 (KERN_ATTR_VECTOR ())
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  u32 w[16];
+
+  w[ 0] = pws[gid].i[ 0];
+  w[ 1] = pws[gid].i[ 1];
+  w[ 2] = pws[gid].i[ 2];
+  w[ 3] = pws[gid].i[ 3];
+  w[ 4] = 0;
+  w[ 5] = 0;
+  w[ 6] = 0;
+  w[ 7] = 0;
+  w[ 8] = 0;
+  w[ 9] = 0;
+  w[10] = 0;
+  w[11] = 0;
+  w[12] = 0;
+  w[13] = 0;
+  w[14] = pws[gid].i[14];
+  w[15] = 0;
+
+  const u32 pw_len = pws[gid].pw_len & 63;
+
+  /**
+   * main
+   */
+
+  m18700s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
+}
+
+__kernel void m18700_s08 (KERN_ATTR_VECTOR ())
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  u32 w[16];
+
+  w[ 0] = pws[gid].i[ 0];
+  w[ 1] = pws[gid].i[ 1];
+  w[ 2] = pws[gid].i[ 2];
+  w[ 3] = pws[gid].i[ 3];
+  w[ 4] = pws[gid].i[ 4];
+  w[ 5] = pws[gid].i[ 5];
+  w[ 6] = pws[gid].i[ 6];
+  w[ 7] = pws[gid].i[ 7];
+  w[ 8] = 0;
+  w[ 9] = 0;
+  w[10] = 0;
+  w[11] = 0;
+  w[12] = 0;
+  w[13] = 0;
+  w[14] = pws[gid].i[14];
+  w[15] = 0;
+
+  const u32 pw_len = pws[gid].pw_len & 63;
+
+  /**
+   * main
+   */
+
+  m18700s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
+}
+
+__kernel void m18700_s16 (KERN_ATTR_VECTOR ())
+{
+  /**
+   * base
+   */
+
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  u32 w[16];
+
+  w[ 0] = pws[gid].i[ 0];
+  w[ 1] = pws[gid].i[ 1];
+  w[ 2] = pws[gid].i[ 2];
+  w[ 3] = pws[gid].i[ 3];
+  w[ 4] = pws[gid].i[ 4];
+  w[ 5] = pws[gid].i[ 5];
+  w[ 6] = pws[gid].i[ 6];
+  w[ 7] = pws[gid].i[ 7];
+  w[ 8] = pws[gid].i[ 8];
+  w[ 9] = pws[gid].i[ 9];
+  w[10] = pws[gid].i[10];
+  w[11] = pws[gid].i[11];
+  w[12] = pws[gid].i[12];
+  w[13] = pws[gid].i[13];
+  w[14] = pws[gid].i[14];
+  w[15] = pws[gid].i[15];
+
+  const u32 pw_len = pws[gid].pw_len & 63;
+
+  /**
+   * main
+   */
+
+  m18700s (w, pw_len, pws, rules_buf, combs_buf, words_buf_r, tmps, hooks, bitmaps_buf_s1_a, bitmaps_buf_s1_b, bitmaps_buf_s1_c, bitmaps_buf_s1_d, bitmaps_buf_s2_a, bitmaps_buf_s2_b, bitmaps_buf_s2_c, bitmaps_buf_s2_d, plains_buf, digests_buf, hashes_shown, salt_bufs, esalt_bufs, d_return_buf, d_extra0_buf, d_extra1_buf, d_extra2_buf, d_extra3_buf, bitmap_mask, bitmap_shift1, bitmap_shift2, salt_pos, loop_pos, loop_cnt, il_cnt, digests_cnt, digests_offset, combs_mode, gid_max);
+}

OpenCL/m18700_a3-pure.cl

@@ -0,0 +1,144 @@
+/**
+ * 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_simd.cl"
+#include "inc_hash_md5.cl"
+
+__kernel void m18700_mxx (KERN_ATTR_VECTOR ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * base
+   */
+
+  const u32 pw_len = pws[gid].pw_len & 255;
+
+  u32x w[64] = { 0 };
+
+  for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
+  {
+    w[idx] = pws[gid].i[idx];
+  }
+
+  /**
+   * loop
+   */
+
+  u32x w0l = w[0];
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
+
+    const u32x w0 = w0l | w0r;
+
+    w[0] = w0;
+
+    u32x hash = 0;
+
+    for (u32 i = 0; i < pw_len; i++)
+    {
+      const u32 c32 = w[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_M_SIMD (r0, r1, r2, r3);
+  }
+}
+
+__kernel void m18700_sxx (KERN_ATTR_VECTOR ())
+{
+  /**
+   * modifier
+   */
+
+  const u64 lid = get_local_id (0);
+  const u64 gid = get_global_id (0);
+
+  if (gid >= gid_max) return;
+
+  /**
+   * digest
+   */
+
+  const u32 search[4] =
+  {
+    digests_buf[digests_offset].digest_buf[DGST_R0],
+    0,
+    0,
+    0
+  };
+
+  /**
+   * base
+   */
+
+  const u32 pw_len = pws[gid].pw_len & 255;
+
+  u32x w[64] = { 0 };
+
+  for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
+  {
+    w[idx] = pws[gid].i[idx];
+  }
+
+  /**
+   * loop
+   */
+
+  u32x w0l = w[0];
+
+  for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
+  {
+    const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
+
+    const u32x w0 = w0l | w0r;
+
+    w[0] = w0;
+
+    u32x hash = 0;
+
+    for (u32 i = 0; i < pw_len; i++)
+    {
+      const u32 c32 = w[i / 4];
+
+      const u32 c = (c32 >> ((i & 3) * 8)) & 0xff;
+
+      hash *= 31;
+      hash += c;
+    }
+
+    const u32x r0 = hash;
+    const u32x r1 = 0;
+    const u32x r2 = 0;
+    const u32x r3 = 0;
+
+    COMPARE_S_SIMD (r0, r1, r2, r3);
+  }
+}