diff --git a/OpenCL/m00600_a0-optimized.cl b/OpenCL/m00600_a0-optimized.cl
index 588f2fcaa..6b7009f7d 100644
--- a/OpenCL/m00600_a0-optimized.cl
+++ b/OpenCL/m00600_a0-optimized.cl
@@ -148,19 +148,19 @@ __kernel void m00600_m04 (__global pw_t *pws, __constant const kernel_rule_t *ru
   u64 tmp_t[2];
   u64 tmp_f[2];
 
-  tmp_h[0] = esalt_bufs->h[0];
-  tmp_h[1] = esalt_bufs->h[1];
-  tmp_h[2] = esalt_bufs->h[2];
-  tmp_h[3] = esalt_bufs->h[3];
-  tmp_h[4] = esalt_bufs->h[4];
-  tmp_h[5] = esalt_bufs->h[5];
-  tmp_h[6] = esalt_bufs->h[6];
-  tmp_h[7] = esalt_bufs->h[7];
+  tmp_h[0] = esalt_bufs[digests_offset].h[0];
+  tmp_h[1] = esalt_bufs[digests_offset].h[1];
+  tmp_h[2] = esalt_bufs[digests_offset].h[2];
+  tmp_h[3] = esalt_bufs[digests_offset].h[3];
+  tmp_h[4] = esalt_bufs[digests_offset].h[4];
+  tmp_h[5] = esalt_bufs[digests_offset].h[5];
+  tmp_h[6] = esalt_bufs[digests_offset].h[6];
+  tmp_h[7] = esalt_bufs[digests_offset].h[7];
 
-  tmp_t[0] = esalt_bufs->t[0];
-  tmp_t[1] = esalt_bufs->t[1];
-  tmp_f[0] = esalt_bufs->f[0];
-  tmp_f[1] = esalt_bufs->f[1];
+  tmp_t[0] = esalt_bufs[digests_offset].t[0];
+  tmp_t[1] = esalt_bufs[digests_offset].t[1];
+  tmp_f[0] = esalt_bufs[digests_offset].f[0];
+  tmp_f[1] = esalt_bufs[digests_offset].f[1];
 
   /**
    * loop
@@ -256,19 +256,19 @@ __kernel void m00600_s04 (__global pw_t *pws, __constant const kernel_rule_t *ru
   u64 tmp_t[2];
   u64 tmp_f[2];
 
-  tmp_h[0] = esalt_bufs->h[0];
-  tmp_h[1] = esalt_bufs->h[1];
-  tmp_h[2] = esalt_bufs->h[2];
-  tmp_h[3] = esalt_bufs->h[3];
-  tmp_h[4] = esalt_bufs->h[4];
-  tmp_h[5] = esalt_bufs->h[5];
-  tmp_h[6] = esalt_bufs->h[6];
-  tmp_h[7] = esalt_bufs->h[7];
+  tmp_h[0] = esalt_bufs[digests_offset].h[0];
+  tmp_h[1] = esalt_bufs[digests_offset].h[1];
+  tmp_h[2] = esalt_bufs[digests_offset].h[2];
+  tmp_h[3] = esalt_bufs[digests_offset].h[3];
+  tmp_h[4] = esalt_bufs[digests_offset].h[4];
+  tmp_h[5] = esalt_bufs[digests_offset].h[5];
+  tmp_h[6] = esalt_bufs[digests_offset].h[6];
+  tmp_h[7] = esalt_bufs[digests_offset].h[7];
 
-  tmp_t[0] = esalt_bufs->t[0];
-  tmp_t[1] = esalt_bufs->t[1];
-  tmp_f[0] = esalt_bufs->f[0];
-  tmp_f[1] = esalt_bufs->f[1];
+  tmp_t[0] = esalt_bufs[digests_offset].t[0];
+  tmp_t[1] = esalt_bufs[digests_offset].t[1];
+  tmp_f[0] = esalt_bufs[digests_offset].f[0];
+  tmp_f[1] = esalt_bufs[digests_offset].f[1];
 
   /**
    * digest
diff --git a/OpenCL/m00600_a1-optimized.cl b/OpenCL/m00600_a1-optimized.cl
index 09b12ac42..829093288 100644
--- a/OpenCL/m00600_a1-optimized.cl
+++ b/OpenCL/m00600_a1-optimized.cl
@@ -148,19 +148,19 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
   u64 tmp_t[2];
   u64 tmp_f[2];
 
-  tmp_h[0] = esalt_bufs->h[0];
-  tmp_h[1] = esalt_bufs->h[1];
-  tmp_h[2] = esalt_bufs->h[2];
-  tmp_h[3] = esalt_bufs->h[3];
-  tmp_h[4] = esalt_bufs->h[4];
-  tmp_h[5] = esalt_bufs->h[5];
-  tmp_h[6] = esalt_bufs->h[6];
-  tmp_h[7] = esalt_bufs->h[7];
+  tmp_h[0] = esalt_bufs[digests_offset].h[0];
+  tmp_h[1] = esalt_bufs[digests_offset].h[1];
+  tmp_h[2] = esalt_bufs[digests_offset].h[2];
+  tmp_h[3] = esalt_bufs[digests_offset].h[3];
+  tmp_h[4] = esalt_bufs[digests_offset].h[4];
+  tmp_h[5] = esalt_bufs[digests_offset].h[5];
+  tmp_h[6] = esalt_bufs[digests_offset].h[6];
+  tmp_h[7] = esalt_bufs[digests_offset].h[7];
 
-  tmp_t[0] = esalt_bufs->t[0];
-  tmp_t[1] = esalt_bufs->t[1];
-  tmp_f[0] = esalt_bufs->f[0];
-  tmp_f[1] = esalt_bufs->f[1];
+  tmp_t[0] = esalt_bufs[digests_offset].t[0];
+  tmp_t[1] = esalt_bufs[digests_offset].t[1];
+  tmp_f[0] = esalt_bufs[digests_offset].f[0];
+  tmp_f[1] = esalt_bufs[digests_offset].f[1];
 
   /**
    * loop
@@ -315,19 +315,19 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   u64 tmp_t[2];
   u64 tmp_f[2];
 
-  tmp_h[0] = esalt_bufs->h[0];
-  tmp_h[1] = esalt_bufs->h[1];
-  tmp_h[2] = esalt_bufs->h[2];
-  tmp_h[3] = esalt_bufs->h[3];
-  tmp_h[4] = esalt_bufs->h[4];
-  tmp_h[5] = esalt_bufs->h[5];
-  tmp_h[6] = esalt_bufs->h[6];
-  tmp_h[7] = esalt_bufs->h[7];
+  tmp_h[0] = esalt_bufs[digests_offset].h[0];
+  tmp_h[1] = esalt_bufs[digests_offset].h[1];
+  tmp_h[2] = esalt_bufs[digests_offset].h[2];
+  tmp_h[3] = esalt_bufs[digests_offset].h[3];
+  tmp_h[4] = esalt_bufs[digests_offset].h[4];
+  tmp_h[5] = esalt_bufs[digests_offset].h[5];
+  tmp_h[6] = esalt_bufs[digests_offset].h[6];
+  tmp_h[7] = esalt_bufs[digests_offset].h[7];
 
-  tmp_t[0] = esalt_bufs->t[0];
-  tmp_t[1] = esalt_bufs->t[1];
-  tmp_f[0] = esalt_bufs->f[0];
-  tmp_f[1] = esalt_bufs->f[1];
+  tmp_t[0] = esalt_bufs[digests_offset].t[0];
+  tmp_t[1] = esalt_bufs[digests_offset].t[1];
+  tmp_f[0] = esalt_bufs[digests_offset].f[0];
+  tmp_f[1] = esalt_bufs[digests_offset].f[1];
 
   /**
    * digest
diff --git a/OpenCL/m00600_a3-optimized.cl b/OpenCL/m00600_a3-optimized.cl
index e172d1a77..ca4da7144 100644
--- a/OpenCL/m00600_a3-optimized.cl
+++ b/OpenCL/m00600_a3-optimized.cl
@@ -132,19 +132,19 @@ __kernel void m00600_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
   u64 tmp_t[2];
   u64 tmp_f[2];
 
-  tmp_h[0] = esalt_bufs->h[0];
-  tmp_h[1] = esalt_bufs->h[1];
-  tmp_h[2] = esalt_bufs->h[2];
-  tmp_h[3] = esalt_bufs->h[3];
-  tmp_h[4] = esalt_bufs->h[4];
-  tmp_h[5] = esalt_bufs->h[5];
-  tmp_h[6] = esalt_bufs->h[6];
-  tmp_h[7] = esalt_bufs->h[7];
+  tmp_h[0] = esalt_bufs[digests_offset].h[0];
+  tmp_h[1] = esalt_bufs[digests_offset].h[1];
+  tmp_h[2] = esalt_bufs[digests_offset].h[2];
+  tmp_h[3] = esalt_bufs[digests_offset].h[3];
+  tmp_h[4] = esalt_bufs[digests_offset].h[4];
+  tmp_h[5] = esalt_bufs[digests_offset].h[5];
+  tmp_h[6] = esalt_bufs[digests_offset].h[6];
+  tmp_h[7] = esalt_bufs[digests_offset].h[7];
 
-  tmp_t[0] = esalt_bufs->t[0];
-  tmp_t[1] = esalt_bufs->t[1];
-  tmp_f[0] = esalt_bufs->f[0];
-  tmp_f[1] = esalt_bufs->f[1];
+  tmp_t[0] = esalt_bufs[digests_offset].t[0];
+  tmp_t[1] = esalt_bufs[digests_offset].t[1];
+  tmp_f[0] = esalt_bufs[digests_offset].f[0];
+  tmp_f[1] = esalt_bufs[digests_offset].f[1];
 
   /**
    * loop
@@ -244,19 +244,19 @@ __kernel void m00600_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
   u64 tmp_t[2];
   u64 tmp_f[2];
 
-  tmp_h[0] = esalt_bufs->h[0];
-  tmp_h[1] = esalt_bufs->h[1];
-  tmp_h[2] = esalt_bufs->h[2];
-  tmp_h[3] = esalt_bufs->h[3];
-  tmp_h[4] = esalt_bufs->h[4];
-  tmp_h[5] = esalt_bufs->h[5];
-  tmp_h[6] = esalt_bufs->h[6];
-  tmp_h[7] = esalt_bufs->h[7];
+  tmp_h[0] = esalt_bufs[digests_offset].h[0];
+  tmp_h[1] = esalt_bufs[digests_offset].h[1];
+  tmp_h[2] = esalt_bufs[digests_offset].h[2];
+  tmp_h[3] = esalt_bufs[digests_offset].h[3];
+  tmp_h[4] = esalt_bufs[digests_offset].h[4];
+  tmp_h[5] = esalt_bufs[digests_offset].h[5];
+  tmp_h[6] = esalt_bufs[digests_offset].h[6];
+  tmp_h[7] = esalt_bufs[digests_offset].h[7];
 
-  tmp_t[0] = esalt_bufs->t[0];
-  tmp_t[1] = esalt_bufs->t[1];
-  tmp_f[0] = esalt_bufs->f[0];
-  tmp_f[1] = esalt_bufs->f[1];
+  tmp_t[0] = esalt_bufs[digests_offset].t[0];
+  tmp_t[1] = esalt_bufs[digests_offset].t[1];
+  tmp_f[0] = esalt_bufs[digests_offset].f[0];
+  tmp_f[1] = esalt_bufs[digests_offset].f[1];
 
   /**
    * digest
diff --git a/docs/changes.txt b/docs/changes.txt
index 8d6d74031..9ba7aab22 100644
--- a/docs/changes.txt
+++ b/docs/changes.txt
@@ -689,6 +689,7 @@ It combines all features of all hashcat projects in one project.
 - Fixed a bug in implementation of DCC2: forced default iteration count for hashes to 10240
 - Fixed a bug in implementation of WPA/WPA2: MAC and nonce stay one their original position as in the hccap file
 - Fixed a bug in implementation of GOST R 34.11-94: zero length passwords were not cracked
+- Fixed a bug in implementation of BLAKE2-512 kernels: incorrect access of the esalt buffer
 
 ##
 ## Technical