diff --git a/OpenCL/m02500.cl b/OpenCL/m02500.cl
index afa6a5776..74ac366a2 100644
--- a/OpenCL/m02500.cl
+++ b/OpenCL/m02500.cl
@@ -41,12 +41,12 @@ void make_kn (u32 *k)
   kr[2] = kr[2] >> 8 | kr[3] << 24;
   kr[3] = kr[3] >> 8;
 
-  kr[3] ^= c * 0x87000000;
-
   k[0] = kl[0] | kr[0];
   k[1] = kl[1] | kr[1];
   k[2] = kl[2] | kr[2];
   k[3] = kl[3] | kr[3];
+
+  k[3] ^= c * 0x87000000;
 }
 
 void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5], u32x digest[5])
diff --git a/OpenCL/m02501.cl b/OpenCL/m02501.cl
index 8a27199e0..b61713868 100644
--- a/OpenCL/m02501.cl
+++ b/OpenCL/m02501.cl
@@ -56,12 +56,12 @@ void make_kn (u32 *k)
   kr[2] = kr[2] >> 8 | kr[3] << 24;
   kr[3] = kr[3] >> 8;
 
-  kr[3] ^= c * 0x87000000;
-
   k[0] = kl[0] | kr[0];
   k[1] = kl[1] | kr[1];
   k[2] = kl[2] | kr[2];
   k[3] = kl[3] | kr[3];
+
+  k[3] ^= c * 0x87000000;
 }
 
 __kernel void m02501_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const wpa_t *wpa_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)