From c15f741dca07f439f324dd348fa3bf4b5356cd78 Mon Sep 17 00:00:00 2001
From: Jens Steube <jens.steube@gmail.com>
Date: Wed, 21 Nov 2018 14:55:22 +0100
Subject: [PATCH] Fixed out-of-boundary read in DPAPI masterkey file v2 OpenCL
 kernel

---
 OpenCL/m15900-pure.cl | 2 ++
 docs/changes.txt      | 1 +
 2 files changed, 3 insertions(+)

diff --git a/OpenCL/m15900-pure.cl b/OpenCL/m15900-pure.cl
index 4a9d8d3a1..e130d6dd4 100644
--- a/OpenCL/m15900-pure.cl
+++ b/OpenCL/m15900-pure.cl
@@ -601,6 +601,8 @@ __kernel void m15900_comp (KERN_ATTR_TMPS_ESALT (dpapimk_tmp_v2_t, dpapimk_t))
     iv[1] = data[1];
     iv[2] = data[2];
     iv[3] = data[3];
+
+    if (contents_off == 32) break;
   }
 
   u32 hmacSalt[4];
diff --git a/docs/changes.txt b/docs/changes.txt
index 877aede7a..e005c4ec5 100644
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -51,6 +51,7 @@
 - Fixed output of --speed-only and --progress-only when fast hashes are used in combination with --slow-candidates
 - Fixed status output of progress value when -S and -l are used in combination
 - Fixed thread count maximum for pure kernels in straight attack mode
+- Fixed out-of-boundary read in DPAPI masterkey file v2 OpenCL kernel
 
 ##
 ## Technical