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 921f064e9..e04b81db6 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;