From 393916c0bfa07b3f703b89225542550cf59dbe23 Mon Sep 17 00:00:00 2001
From: "R. Yushaev" <44146334+Naufragous@users.noreply.github.com>
Date: Tue, 18 Dec 2018 14:32:56 +0100
Subject: [PATCH] Allow cracking non-unique salts for Office 2013

With hash-mode 9600 (MS Office 2013) there can be multiple hashes with
the same salt but with different encryption verifiers in esalt_bufs.
This commit adds the functionality to execute _comp kernels for
different hashes after deriving their common key once.

Fixes #1826
---
 OpenCL/m09600-pure.cl | 10 ++++++----
 src/opencl.c          | 16 ++++++++++++++++
 2 files changed, 22 insertions(+), 4 deletions(-)

diff --git a/OpenCL/m09600-pure.cl b/OpenCL/m09600-pure.cl
index 4f59fa484..27a01a584 100644
--- a/OpenCL/m09600-pure.cl
+++ b/OpenCL/m09600-pure.cl
@@ -354,12 +354,14 @@ __kernel void m09600_comp (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t)
 
   AES256_set_decrypt_key (ks, ukey, s_te0, s_te1, s_te2, s_te3, s_te4, s_td0, s_td1, s_td2, s_td3, s_td4);
 
+  const u32 digest_cur = digests_offset + loop_pos;
+
   u32 data[4];
 
-  data[0] = esalt_bufs[digests_offset].encryptedVerifier[0];
-  data[1] = esalt_bufs[digests_offset].encryptedVerifier[1];
-  data[2] = esalt_bufs[digests_offset].encryptedVerifier[2];
-  data[3] = esalt_bufs[digests_offset].encryptedVerifier[3];
+  data[0] = esalt_bufs[digest_cur].encryptedVerifier[0];
+  data[1] = esalt_bufs[digest_cur].encryptedVerifier[1];
+  data[2] = esalt_bufs[digest_cur].encryptedVerifier[2];
+  data[3] = esalt_bufs[digest_cur].encryptedVerifier[3];
 
   u32 out[4];
 
diff --git a/src/opencl.c b/src/opencl.c
index ecb824d20..5e87a54ba 100644
--- a/src/opencl.c
+++ b/src/opencl.c
@@ -1481,6 +1481,22 @@ int choose_kernel (hashcat_ctx_t *hashcat_ctx, hc_device_param_t *device_param,
           if (status_ctx->run_thread_level2 == false) break;
         }
       }
+      else if (hashconfig->hash_mode == 9600)
+      {
+        const u32 loops_cnt = hashes->salts_buf[salt_pos].digests_cnt;
+
+        for (u32 loops_pos = 0; loops_pos < loops_cnt; loops_pos++)
+        {
+          device_param->kernel_params_buf32[28] = loops_pos;
+          device_param->kernel_params_buf32[29] = loops_cnt;
+
+          CL_rc = run_kernel (hashcat_ctx, device_param, KERN_RUN_3, pws_cnt, false, 0);
+
+          if (CL_rc == -1) return -1;
+
+          if (status_ctx->run_thread_level2 == false) break;
+        }
+      }
       else if ((hashconfig->hash_mode == 16800) || (hashconfig->hash_mode == 16801))
       {
         const u32 loops_cnt = hashes->salts_buf[salt_pos].digests_cnt;