Increase 1Password, cloudkeychain cracking performance

diff --git a/OpenCL/inc_vendor.cl b/OpenCL/inc_vendor.cl
index 3adde96..87a1d08 100644 (file)
--- a/OpenCL/inc_vendor.cl
+++ b/OpenCL/inc_vendor.cl
@@ -101,9 +101,6 @@
 #if KERN_TYPE == 7900
 #undef _unroll
 #endif
-#if KERN_TYPE == 8200
-#undef _unroll
-#endif
 #if KERN_TYPE == 10400
 #undef _unroll
 #endif

diff --git a/OpenCL/m08200.cl b/OpenCL/m08200.cl
index ae7e6de..cf1bd15 100644 (file)
--- a/OpenCL/m08200.cl
+++ b/OpenCL/m08200.cl
@@ -385,6 +385,59 @@ void hmac_sha512_run (const u64 w1[16], const u64 ipad[8], const u64 opad[8], u6
   sha512_transform (w, dgst);
 }
 
+void hmac_sha512_run_x (const u64 ipad[8], const u64 opad[8], u64 dgst[8])
+{
+  u64 w[16];
+
+  w[ 0] = dgst[0];
+  w[ 1] = dgst[1];
+  w[ 2] = dgst[2];
+  w[ 3] = dgst[3];
+  w[ 4] = dgst[4];
+  w[ 5] = dgst[5];
+  w[ 6] = dgst[6];
+  w[ 7] = dgst[7];
+  w[ 8] = 0x8000000000000000;
+  w[ 9] = 0;
+  w[10] = 0;
+  w[11] = 0;
+  w[12] = 0;
+  w[13] = 0;
+  w[14] = 0;
+  w[15] = (128 + 64) * 8;
+
+  dgst[0] = ipad[0];
+  dgst[1] = ipad[1];
+  dgst[2] = ipad[2];
+  dgst[3] = ipad[3];
+  dgst[4] = ipad[4];
+  dgst[5] = ipad[5];
+  dgst[6] = ipad[6];
+  dgst[7] = ipad[7];
+
+  sha512_transform (w, dgst);
+
+  w[ 0] = dgst[0];
+  w[ 1] = dgst[1];
+  w[ 2] = dgst[2];
+  w[ 3] = dgst[3];
+  w[ 4] = dgst[4];
+  w[ 5] = dgst[5];
+  w[ 6] = dgst[6];
+  w[ 7] = dgst[7];
+
+  dgst[0] = opad[0];
+  dgst[1] = opad[1];
+  dgst[2] = opad[2];
+  dgst[3] = opad[3];
+  dgst[4] = opad[4];
+  dgst[5] = opad[5];
+  dgst[6] = opad[6];
+  dgst[7] = opad[7];
+
+  sha512_transform (w, dgst);
+}
+
 void hmac_sha512_init (u64 w[16], u64 ipad[8], u64 opad[8])
 {
   w[ 0] ^= 0x3636363636363636;
@@ -617,81 +670,39 @@ __kernel void m08200_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf
   opad[6] = tmps[gid].opad[6];
   opad[7] = tmps[gid].opad[7];
 
-  for (u32 i = 0; i < 8; i += 8)
-  {
-    u64 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];
-
-    u64 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++)
-    {
-      u64 w[16];
-
-      w[ 0] = dgst[0];
-      w[ 1] = dgst[1];
-      w[ 2] = dgst[2];
-      w[ 3] = dgst[3];
-      w[ 4] = dgst[4];
-      w[ 5] = dgst[5];
-      w[ 6] = dgst[6];
-      w[ 7] = dgst[7];
-      w[ 8] = 0x8000000000000000;
-      w[ 9] = 0;
-      w[10] = 0;
-      w[11] = 0;
-      w[12] = 0;
-      w[13] = 0;
-      w[14] = 0;
-      w[15] = (128 + 64) * 8;
-
-      hmac_sha512_run (w, 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];
-    }
+  u64 dgst[8];
 
-    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];
+  dgst[0] = tmps[gid].dgst[0];
+  dgst[1] = tmps[gid].dgst[1];
+  dgst[2] = tmps[gid].dgst[2];
+  dgst[3] = tmps[gid].dgst[3];
+  dgst[4] = tmps[gid].dgst[4];
+  dgst[5] = tmps[gid].dgst[5];
+  dgst[6] = tmps[gid].dgst[6];
+  dgst[7] = tmps[gid].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];
+  for (u32 j = 0; j < loop_cnt; j++)
+  {
+    hmac_sha512_run_x (ipad, opad, dgst);
+
+    tmps[gid].out[0] ^= dgst[0];
+    tmps[gid].out[1] ^= dgst[1];
+    tmps[gid].out[2] ^= dgst[2];
+    tmps[gid].out[3] ^= dgst[3];
+    tmps[gid].out[4] ^= dgst[4];
+    tmps[gid].out[5] ^= dgst[5];
+    tmps[gid].out[6] ^= dgst[6];
+    tmps[gid].out[7] ^= dgst[7];
   }
+
+  tmps[gid].dgst[0] = dgst[0];
+  tmps[gid].dgst[1] = dgst[1];
+  tmps[gid].dgst[2] = dgst[2];
+  tmps[gid].dgst[3] = dgst[3];
+  tmps[gid].dgst[4] = dgst[4];
+  tmps[gid].dgst[5] = dgst[5];
+  tmps[gid].dgst[6] = dgst[6];
+  tmps[gid].dgst[7] = dgst[7];
 }
 
 __kernel void m08200_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global pbkdf2_sha512_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 cloudkey_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)