Increase Oracle T: Type (Oracle 12+) cracking performance

diff --git a/OpenCL/m12300.cl b/OpenCL/m12300.cl
index 3a87faa..a66e086 100644 (file)
--- a/OpenCL/m12300.cl
+++ b/OpenCL/m12300.cl
@@ -176,6 +176,59 @@ void hmac_run (const u64 w1[16], const u64 ipad[8], const u64 opad[8], u64 dgst[
   sha512_transform (w, dgst);
 }
 
+void hmac_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_init (u64 w[16], u64 ipad[8], u64 opad[8])
 {
   w[ 0] ^= 0x3636363636363636;
@@ -392,81 +445,39 @@ __kernel void m12300_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_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_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 m12300_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global oraclet_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 void *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)