diff --git a/OpenCL/m10800_a0.cl b/OpenCL/m10800_a0.cl
index abe6efa67..6710d0628 100644
--- a/OpenCL/m10800_a0.cl
+++ b/OpenCL/m10800_a0.cl
@@ -30,14 +30,7 @@ __kernel void m10800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len;
-
-  u32 w[64] = { 0 };
-
-  for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
-  {
-    w[idx] = swap32_S (pws[gid].i[idx]);
-  }
+  COPY_PW (pws[gid]);
 
   /**
    * loop
@@ -53,7 +46,7 @@ __kernel void m10800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru
 
     sha384_init (&ctx);
 
-    sha384_update (&ctx, tmp.i, tmp.pw_len);
+    sha384_update_swap (&ctx, tmp.i, tmp.pw_len);
 
     sha384_final (&ctx);
 
@@ -93,14 +86,7 @@ __kernel void m10800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru
    * base
    */
 
-  const u32 pw_len = pws[gid].pw_len;
-
-  u32 w[64] = { 0 };
-
-  for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
-  {
-    w[idx] = swap32_S (pws[gid].i[idx]);
-  }
+  COPY_PW (pws[gid]);
 
   /**
    * loop
@@ -116,7 +102,7 @@ __kernel void m10800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru
 
     sha384_init (&ctx);
 
-    sha384_update (&ctx, tmp.i, tmp.pw_len);
+    sha384_update_swap (&ctx, tmp.i, tmp.pw_len);
 
     sha384_final (&ctx);