diff --git a/OpenCL/m00670_a0.cl b/OpenCL/m00670_a0.cl
index 2d9c78d0a..ffb0fecb4 100644
--- a/OpenCL/m00670_a0.cl
+++ b/OpenCL/m00670_a0.cl
@@ -132,6 +132,17 @@ __kernel void m00670_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
 
   const u32 pw_len = pws[gid].pw_len;
 
+  u32 iv[2]    = { 0 };
+  u32 plain[2] = { 0 };
+  u32 position = esalt_bufs->position;
+  u32 plain_length = esalt_bufs->plain_length;
+
+  iv[0] = esalt_bufs->iv[0];
+  iv[1] = esalt_bufs->iv[1];
+
+  plain[0] = esalt_bufs->plain[0];
+  plain[1] = esalt_bufs->plain[1];
+
   /**
    * loop
    */
@@ -145,13 +156,14 @@ __kernel void m00670_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
 
     const u32x out_len = apply_rules_vect(pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1);
 
-    u64x digest[8];
+    u32x digest[4] = { 0 };
 
+    chacha20_transform (w0, w1, position, iv, plain, digest);
 
-    const u32x r0 = h32_from_64(digest[0]);
-    const u32x r1 = l32_from_64(digest[0]);
-    const u32x r2 = h32_from_64(digest[1]);
-    const u32x r3 = l32_from_64(digest[1]);
+    const u32x r0 = digest[0];
+    const u32x r1 = digest[1];
+    const u32x r2 = digest[2];
+    const u32x r3 = digest[3];
 
     COMPARE_M_SIMD(r0, r1, r2, r3);
   }