diff --git a/OpenCL/m00000_a1-pure.cl b/OpenCL/m00000_a1-pure.cl
index fef760853..8d9a69b8f 100644
--- a/OpenCL/m00000_a1-pure.cl
+++ b/OpenCL/m00000_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m00000_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m00000_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
diff --git a/OpenCL/m00010_a1-pure.cl b/OpenCL/m00010_a1-pure.cl
index 8ae461f10..fbcc5df3b 100644
--- a/OpenCL/m00010_a1-pure.cl
+++ b/OpenCL/m00010_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m00010_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m00010_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m00020_a1-pure.cl b/OpenCL/m00020_a1-pure.cl
index e862c48eb..dcb9281ae 100644
--- a/OpenCL/m00020_a1-pure.cl
+++ b/OpenCL/m00020_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m00020_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m00020_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
diff --git a/OpenCL/m00030_a1-pure.cl b/OpenCL/m00030_a1-pure.cl
index 83ed6766a..5192b9660 100644
--- a/OpenCL/m00030_a1-pure.cl
+++ b/OpenCL/m00030_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m00030_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m00030_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m00040_a1-pure.cl b/OpenCL/m00040_a1-pure.cl
index 67fa78329..163c5dc52 100644
--- a/OpenCL/m00040_a1-pure.cl
+++ b/OpenCL/m00040_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m00040_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m00040_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
diff --git a/OpenCL/m00100_a1-pure.cl b/OpenCL/m00100_a1-pure.cl
index 375bedbd7..fb29cf3f0 100644
--- a/OpenCL/m00100_a1-pure.cl
+++ b/OpenCL/m00100_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m00100_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m00100_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
diff --git a/OpenCL/m00110_a1-pure.cl b/OpenCL/m00110_a1-pure.cl
index 2bc8177ca..9ff685d3e 100644
--- a/OpenCL/m00110_a1-pure.cl
+++ b/OpenCL/m00110_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m00110_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m00110_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m00120_a1-pure.cl b/OpenCL/m00120_a1-pure.cl
index 8acbd83b4..ef0ff1cc4 100644
--- a/OpenCL/m00120_a1-pure.cl
+++ b/OpenCL/m00120_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m00120_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m00120_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
diff --git a/OpenCL/m00130_a1-pure.cl b/OpenCL/m00130_a1-pure.cl
index a9b17120c..2d254c7b2 100644
--- a/OpenCL/m00130_a1-pure.cl
+++ b/OpenCL/m00130_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m00130_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m00130_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m00140_a1-pure.cl b/OpenCL/m00140_a1-pure.cl
index 02842ffc0..ae7cdf2a1 100644
--- a/OpenCL/m00140_a1-pure.cl
+++ b/OpenCL/m00140_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m00140_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m00140_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
diff --git a/OpenCL/m00300_a1-pure.cl b/OpenCL/m00300_a1-pure.cl
index 67edd34b8..cdd334522 100644
--- a/OpenCL/m00300_a1-pure.cl
+++ b/OpenCL/m00300_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m00300_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx0;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
@@ -119,7 +119,7 @@ __kernel void m00300_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx0;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
diff --git a/OpenCL/m00900_a1-pure.cl b/OpenCL/m00900_a1-pure.cl
index 50118c203..22763af94 100644
--- a/OpenCL/m00900_a1-pure.cl
+++ b/OpenCL/m00900_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m00900_mxx (KERN_ATTR_BASIC ())
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m00900_sxx (KERN_ATTR_BASIC ())
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
diff --git a/OpenCL/m01000_a1-pure.cl b/OpenCL/m01000_a1-pure.cl
index c5b438e0b..264e57376 100644
--- a/OpenCL/m01000_a1-pure.cl
+++ b/OpenCL/m01000_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m01000_mxx (KERN_ATTR_BASIC ())
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m01000_sxx (KERN_ATTR_BASIC ())
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
diff --git a/OpenCL/m01100_a1-pure.cl b/OpenCL/m01100_a1-pure.cl
index a6f8b7c81..ddfdda6e9 100644
--- a/OpenCL/m01100_a1-pure.cl
+++ b/OpenCL/m01100_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m01100_mxx (KERN_ATTR_BASIC ())
   {
     md4_ctx_t ctx1 = ctx0;
 
-    md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx1);
 
@@ -129,7 +129,7 @@ __kernel void m01100_sxx (KERN_ATTR_BASIC ())
   {
     md4_ctx_t ctx1 = ctx0;
 
-    md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx1);
 
diff --git a/OpenCL/m01300_a1-pure.cl b/OpenCL/m01300_a1-pure.cl
index fcb155729..d71b75297 100644
--- a/OpenCL/m01300_a1-pure.cl
+++ b/OpenCL/m01300_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m01300_mxx (KERN_ATTR_BASIC ())
   {
     sha224_ctx_t ctx = ctx0;
 
-    sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha224_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m01300_sxx (KERN_ATTR_BASIC ())
   {
     sha224_ctx_t ctx = ctx0;
 
-    sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha224_final (&ctx);
 
diff --git a/OpenCL/m01400_a1-pure.cl b/OpenCL/m01400_a1-pure.cl
index fc040002b..1518aaf2b 100644
--- a/OpenCL/m01400_a1-pure.cl
+++ b/OpenCL/m01400_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m01400_mxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m01400_sxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_final (&ctx);
 
diff --git a/OpenCL/m01410_a1-pure.cl b/OpenCL/m01410_a1-pure.cl
index eaa6cbae0..d6759f85a 100644
--- a/OpenCL/m01410_a1-pure.cl
+++ b/OpenCL/m01410_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m01410_mxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m01410_sxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m01420_a1-pure.cl b/OpenCL/m01420_a1-pure.cl
index 429e62e2e..60bcfe246 100644
--- a/OpenCL/m01420_a1-pure.cl
+++ b/OpenCL/m01420_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m01420_mxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m01420_sxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_final (&ctx);
 
diff --git a/OpenCL/m01430_a1-pure.cl b/OpenCL/m01430_a1-pure.cl
index 1f441e83c..3c3554d9e 100644
--- a/OpenCL/m01430_a1-pure.cl
+++ b/OpenCL/m01430_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m01430_mxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m01430_sxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m01440_a1-pure.cl b/OpenCL/m01440_a1-pure.cl
index c3e7070b4..2f4429df9 100644
--- a/OpenCL/m01440_a1-pure.cl
+++ b/OpenCL/m01440_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m01440_mxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m01440_sxx (KERN_ATTR_BASIC ())
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_final (&ctx);
 
diff --git a/OpenCL/m01700_a1-pure.cl b/OpenCL/m01700_a1-pure.cl
index 038558998..ba873edfb 100644
--- a/OpenCL/m01700_a1-pure.cl
+++ b/OpenCL/m01700_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m01700_mxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m01700_sxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_final (&ctx);
 
diff --git a/OpenCL/m01710_a1-pure.cl b/OpenCL/m01710_a1-pure.cl
index 296a74027..a3c4ffa82 100644
--- a/OpenCL/m01710_a1-pure.cl
+++ b/OpenCL/m01710_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m01710_mxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m01710_sxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m01720_a1-pure.cl b/OpenCL/m01720_a1-pure.cl
index 3fa6901f4..fdb7660a7 100644
--- a/OpenCL/m01720_a1-pure.cl
+++ b/OpenCL/m01720_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m01720_mxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m01720_sxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_final (&ctx);
 
diff --git a/OpenCL/m01730_a1-pure.cl b/OpenCL/m01730_a1-pure.cl
index 77c856527..2e60d7599 100644
--- a/OpenCL/m01730_a1-pure.cl
+++ b/OpenCL/m01730_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m01730_mxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m01730_sxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m01740_a1-pure.cl b/OpenCL/m01740_a1-pure.cl
index 0f0d9692b..3152b02d6 100644
--- a/OpenCL/m01740_a1-pure.cl
+++ b/OpenCL/m01740_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m01740_mxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m01740_sxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_final (&ctx);
 
diff --git a/OpenCL/m02610_a1-pure.cl b/OpenCL/m02610_a1-pure.cl
index a36061bf5..7481d46ef 100644
--- a/OpenCL/m02610_a1-pure.cl
+++ b/OpenCL/m02610_a1-pure.cl
@@ -81,7 +81,7 @@ __kernel void m02610_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
@@ -194,7 +194,7 @@ __kernel void m02610_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
diff --git a/OpenCL/m02810_a1-pure.cl b/OpenCL/m02810_a1-pure.cl
index 17dc88fbe..eea4f2949 100644
--- a/OpenCL/m02810_a1-pure.cl
+++ b/OpenCL/m02810_a1-pure.cl
@@ -81,7 +81,7 @@ __kernel void m02810_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
@@ -217,7 +217,7 @@ __kernel void m02810_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
diff --git a/OpenCL/m03710_a1-pure.cl b/OpenCL/m03710_a1-pure.cl
index 5455decfa..52fa8d01a 100644
--- a/OpenCL/m03710_a1-pure.cl
+++ b/OpenCL/m03710_a1-pure.cl
@@ -81,7 +81,7 @@ __kernel void m03710_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
@@ -207,7 +207,7 @@ __kernel void m03710_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
diff --git a/OpenCL/m03800_a1-pure.cl b/OpenCL/m03800_a1-pure.cl
index fecfd6556..b78e41a2f 100644
--- a/OpenCL/m03800_a1-pure.cl
+++ b/OpenCL/m03800_a1-pure.cl
@@ -53,7 +53,7 @@ __kernel void m03800_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx, s, salt_len);
 
@@ -120,7 +120,7 @@ __kernel void m03800_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m03910_a1-pure.cl b/OpenCL/m03910_a1-pure.cl
index b5f36d7a2..cf947820f 100644
--- a/OpenCL/m03910_a1-pure.cl
+++ b/OpenCL/m03910_a1-pure.cl
@@ -81,7 +81,7 @@ __kernel void m03910_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
@@ -217,7 +217,7 @@ __kernel void m03910_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
diff --git a/OpenCL/m04010_a1-pure.cl b/OpenCL/m04010_a1-pure.cl
index 7a8998795..d14fd1a2a 100644
--- a/OpenCL/m04010_a1-pure.cl
+++ b/OpenCL/m04010_a1-pure.cl
@@ -76,7 +76,7 @@ __kernel void m04010_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0t;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
@@ -193,7 +193,7 @@ __kernel void m04010_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0t;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
diff --git a/OpenCL/m04110_a1-pure.cl b/OpenCL/m04110_a1-pure.cl
index 1f3c9af3a..07be76686 100644
--- a/OpenCL/m04110_a1-pure.cl
+++ b/OpenCL/m04110_a1-pure.cl
@@ -87,7 +87,7 @@ __kernel void m04110_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0t;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx1, s, salt_len);
 
@@ -217,7 +217,7 @@ __kernel void m04110_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0t;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx1, s, salt_len);
 
diff --git a/OpenCL/m04310_a1-pure.cl b/OpenCL/m04310_a1-pure.cl
index 37ba88f15..3600a5665 100644
--- a/OpenCL/m04310_a1-pure.cl
+++ b/OpenCL/m04310_a1-pure.cl
@@ -81,7 +81,7 @@ __kernel void m04310_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
@@ -194,7 +194,7 @@ __kernel void m04310_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
diff --git a/OpenCL/m04400_a1-pure.cl b/OpenCL/m04400_a1-pure.cl
index 4841e5398..e81f43c0a 100644
--- a/OpenCL/m04400_a1-pure.cl
+++ b/OpenCL/m04400_a1-pure.cl
@@ -73,7 +73,7 @@ __kernel void m04400_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx0;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
@@ -180,7 +180,7 @@ __kernel void m04400_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx0;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
diff --git a/OpenCL/m04500_a1-pure.cl b/OpenCL/m04500_a1-pure.cl
index fe9ce3ce3..774b2f937 100644
--- a/OpenCL/m04500_a1-pure.cl
+++ b/OpenCL/m04500_a1-pure.cl
@@ -72,7 +72,7 @@ __kernel void m04500_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx0;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
@@ -179,7 +179,7 @@ __kernel void m04500_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx0;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
diff --git a/OpenCL/m04520_a1-pure.cl b/OpenCL/m04520_a1-pure.cl
index b5488cbca..a89957884 100644
--- a/OpenCL/m04520_a1-pure.cl
+++ b/OpenCL/m04520_a1-pure.cl
@@ -78,7 +78,7 @@ __kernel void m04520_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx1l;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
@@ -200,7 +200,7 @@ __kernel void m04520_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx1l;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
diff --git a/OpenCL/m04700_a1-pure.cl b/OpenCL/m04700_a1-pure.cl
index 7da0e4f5f..fff7fa469 100644
--- a/OpenCL/m04700_a1-pure.cl
+++ b/OpenCL/m04700_a1-pure.cl
@@ -73,7 +73,7 @@ __kernel void m04700_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
@@ -175,7 +175,7 @@ __kernel void m04700_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
diff --git a/OpenCL/m04800_a1-pure.cl b/OpenCL/m04800_a1-pure.cl
index 6e527c17f..bacfb956b 100644
--- a/OpenCL/m04800_a1-pure.cl
+++ b/OpenCL/m04800_a1-pure.cl
@@ -55,7 +55,7 @@ __kernel void m04800_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx, s, salt_len);
 
@@ -124,7 +124,7 @@ __kernel void m04800_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m04900_a1-pure.cl b/OpenCL/m04900_a1-pure.cl
index 17c62a3ea..dc873e1ce 100644
--- a/OpenCL/m04900_a1-pure.cl
+++ b/OpenCL/m04900_a1-pure.cl
@@ -53,7 +53,7 @@ __kernel void m04900_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx, s, salt_len);
 
@@ -120,7 +120,7 @@ __kernel void m04900_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m05100_a1-pure.cl b/OpenCL/m05100_a1-pure.cl
index ceac4ee19..646164b03 100644
--- a/OpenCL/m05100_a1-pure.cl
+++ b/OpenCL/m05100_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m05100_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m05100_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
diff --git a/OpenCL/m05500_a1-pure.cl b/OpenCL/m05500_a1-pure.cl
index a09d65d30..8011f935e 100644
--- a/OpenCL/m05500_a1-pure.cl
+++ b/OpenCL/m05500_a1-pure.cl
@@ -555,7 +555,7 @@ __kernel void m05500_mxx (KERN_ATTR_BASIC ())
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
@@ -689,7 +689,7 @@ __kernel void m05500_sxx (KERN_ATTR_BASIC ())
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
diff --git a/OpenCL/m05600_a1-pure.cl b/OpenCL/m05600_a1-pure.cl
index 581389943..8db775138 100644
--- a/OpenCL/m05600_a1-pure.cl
+++ b/OpenCL/m05600_a1-pure.cl
@@ -43,7 +43,7 @@ __kernel void m05600_mxx (KERN_ATTR_ESALT (netntlm_t))
   {
     md4_ctx_t ctx1 = ctx10;
 
-    md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx1);
 
@@ -152,7 +152,7 @@ __kernel void m05600_sxx (KERN_ATTR_ESALT (netntlm_t))
   {
     md4_ctx_t ctx1 = ctx10;
 
-    md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx1);
 
diff --git a/OpenCL/m06000_a1-pure.cl b/OpenCL/m06000_a1-pure.cl
index 4fa59fe25..321bbc337 100644
--- a/OpenCL/m06000_a1-pure.cl
+++ b/OpenCL/m06000_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m06000_mxx (KERN_ATTR_BASIC ())
   {
     ripemd160_ctx_t ctx = ctx0;
 
-    ripemd160_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    ripemd160_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     ripemd160_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m06000_sxx (KERN_ATTR_BASIC ())
   {
     ripemd160_ctx_t ctx = ctx0;
 
-    ripemd160_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    ripemd160_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     ripemd160_final (&ctx);
 
diff --git a/OpenCL/m06100_a1-pure.cl b/OpenCL/m06100_a1-pure.cl
index 16cbec86b..554887732 100644
--- a/OpenCL/m06100_a1-pure.cl
+++ b/OpenCL/m06100_a1-pure.cl
@@ -73,7 +73,7 @@ __kernel void m06100_mxx (KERN_ATTR_BASIC ())
   {
     whirlpool_ctx_t ctx = ctx0;
 
-    whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     whirlpool_final (&ctx);
 
@@ -158,7 +158,7 @@ __kernel void m06100_sxx (KERN_ATTR_BASIC ())
   {
     whirlpool_ctx_t ctx = ctx0;
 
-    whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     whirlpool_final (&ctx);
 
diff --git a/OpenCL/m07000_a1-pure.cl b/OpenCL/m07000_a1-pure.cl
index 521b13d63..496340b8a 100644
--- a/OpenCL/m07000_a1-pure.cl
+++ b/OpenCL/m07000_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m07000_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     /**
      * pepper
@@ -128,7 +128,7 @@ __kernel void m07000_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     /**
      * pepper
diff --git a/OpenCL/m07500_a1-pure.cl b/OpenCL/m07500_a1-pure.cl
index c80345831..5102d11aa 100644
--- a/OpenCL/m07500_a1-pure.cl
+++ b/OpenCL/m07500_a1-pure.cl
@@ -317,7 +317,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_A
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
@@ -386,7 +386,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_sxx (KERN_A
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
diff --git a/OpenCL/m08300_a1-pure.cl b/OpenCL/m08300_a1-pure.cl
index 27a24da49..54ce3028e 100644
--- a/OpenCL/m08300_a1-pure.cl
+++ b/OpenCL/m08300_a1-pure.cl
@@ -64,7 +64,7 @@ __kernel void m08300_mxx (KERN_ATTR_BASIC ())
 
     sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len & 255);
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx1, s_pc, salt_len_pc + 1);
 
@@ -179,7 +179,7 @@ __kernel void m08300_sxx (KERN_ATTR_BASIC ())
 
     sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len & 255);
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx1, s_pc, salt_len_pc + 1);
 
diff --git a/OpenCL/m08400_a1-pure.cl b/OpenCL/m08400_a1-pure.cl
index 6a80b86dd..e309aeea4 100644
--- a/OpenCL/m08400_a1-pure.cl
+++ b/OpenCL/m08400_a1-pure.cl
@@ -78,7 +78,7 @@ __kernel void m08400_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx1l;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
@@ -239,7 +239,7 @@ __kernel void m08400_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx1l;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
diff --git a/OpenCL/m10800_a1-pure.cl b/OpenCL/m10800_a1-pure.cl
index 267ffac25..fec24179c 100644
--- a/OpenCL/m10800_a1-pure.cl
+++ b/OpenCL/m10800_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m10800_mxx (KERN_ATTR_BASIC ())
   {
     sha384_ctx_t ctx = ctx0;
 
-    sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha384_final (&ctx);
 
@@ -96,7 +96,7 @@ __kernel void m10800_sxx (KERN_ATTR_BASIC ())
   {
     sha384_ctx_t ctx = ctx0;
 
-    sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha384_final (&ctx);
 
diff --git a/OpenCL/m11000_a1-pure.cl b/OpenCL/m11000_a1-pure.cl
index 38a9cce36..f3ea91ed9 100644
--- a/OpenCL/m11000_a1-pure.cl
+++ b/OpenCL/m11000_a1-pure.cl
@@ -44,7 +44,7 @@ __kernel void m11000_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
@@ -100,7 +100,7 @@ __kernel void m11000_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx);
 
diff --git a/OpenCL/m11100_a1-pure.cl b/OpenCL/m11100_a1-pure.cl
index 6046dd6e7..c82ce5603 100644
--- a/OpenCL/m11100_a1-pure.cl
+++ b/OpenCL/m11100_a1-pure.cl
@@ -94,7 +94,7 @@ __kernel void m11100_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0t;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     u32 s0[4];
     u32 s1[4];
@@ -250,7 +250,7 @@ __kernel void m11100_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx1 = ctx0t;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     u32 s0[4];
     u32 s1[4];
diff --git a/OpenCL/m11200_a1-pure.cl b/OpenCL/m11200_a1-pure.cl
index eedfe51ea..2f482eaf2 100644
--- a/OpenCL/m11200_a1-pure.cl
+++ b/OpenCL/m11200_a1-pure.cl
@@ -48,7 +48,7 @@ __kernel void m11200_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx2 = ctx2l;
 
-    sha1_update_global_swap (&ctx2, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx2, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx2);
 
@@ -174,7 +174,7 @@ __kernel void m11200_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx2 = ctx2l;
 
-    sha1_update_global_swap (&ctx2, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx2, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx2);
 
diff --git a/OpenCL/m11400_a1-pure.cl b/OpenCL/m11400_a1-pure.cl
index 7067a3325..78b12420c 100644
--- a/OpenCL/m11400_a1-pure.cl
+++ b/OpenCL/m11400_a1-pure.cl
@@ -74,7 +74,7 @@ __kernel void m11400_mxx (KERN_ATTR_ESALT (sip_t))
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
@@ -180,7 +180,7 @@ __kernel void m11400_sxx (KERN_ATTR_ESALT (sip_t))
   {
     md5_ctx_t ctx1 = ctx0;
 
-    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md5_final (&ctx1);
 
diff --git a/OpenCL/m11700_a1-pure.cl b/OpenCL/m11700_a1-pure.cl
index 8b05a6969..81b1c6d25 100644
--- a/OpenCL/m11700_a1-pure.cl
+++ b/OpenCL/m11700_a1-pure.cl
@@ -71,7 +71,7 @@ __kernel void m11700_mxx (KERN_ATTR_BASIC ())
   {
     streebog256_ctx_t ctx = ctx0;
 
-    streebog256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    streebog256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     streebog256_final (&ctx);
 
@@ -154,7 +154,7 @@ __kernel void m11700_sxx (KERN_ATTR_BASIC ())
   {
     streebog256_ctx_t ctx = ctx0;
 
-    streebog256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    streebog256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     streebog256_final (&ctx);
 
diff --git a/OpenCL/m11800_a1-pure.cl b/OpenCL/m11800_a1-pure.cl
index e4c63e3f9..1e25cffce 100644
--- a/OpenCL/m11800_a1-pure.cl
+++ b/OpenCL/m11800_a1-pure.cl
@@ -71,7 +71,7 @@ __kernel void m11800_mxx (KERN_ATTR_BASIC ())
   {
     streebog512_ctx_t ctx = ctx0;
 
-    streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     streebog512_final (&ctx);
 
@@ -154,7 +154,7 @@ __kernel void m11800_sxx (KERN_ATTR_BASIC ())
   {
     streebog512_ctx_t ctx = ctx0;
 
-    streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     streebog512_final (&ctx);
 
diff --git a/OpenCL/m12600_a1-pure.cl b/OpenCL/m12600_a1-pure.cl
index c0940de6e..95d6b4fcc 100644
--- a/OpenCL/m12600_a1-pure.cl
+++ b/OpenCL/m12600_a1-pure.cl
@@ -88,7 +88,7 @@ __kernel void m12600_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx0;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
@@ -234,7 +234,7 @@ __kernel void m12600_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx0;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
diff --git a/OpenCL/m13100_a1-pure.cl b/OpenCL/m13100_a1-pure.cl
index 26166dbb9..ce6ad8db0 100644
--- a/OpenCL/m13100_a1-pure.cl
+++ b/OpenCL/m13100_a1-pure.cl
@@ -415,7 +415,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_A
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
@@ -475,7 +475,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_sxx (KERN_A
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
diff --git a/OpenCL/m13300_a1-pure.cl b/OpenCL/m13300_a1-pure.cl
index 6ce60d965..917239f57 100644
--- a/OpenCL/m13300_a1-pure.cl
+++ b/OpenCL/m13300_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m13300_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
@@ -98,7 +98,7 @@ __kernel void m13300_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
diff --git a/OpenCL/m13500_a1-pure.cl b/OpenCL/m13500_a1-pure.cl
index 37eb59adb..506bf0482 100644
--- a/OpenCL/m13500_a1-pure.cl
+++ b/OpenCL/m13500_a1-pure.cl
@@ -71,7 +71,7 @@ __kernel void m13500_mxx (KERN_ATTR_ESALT (pstoken_t))
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
@@ -154,7 +154,7 @@ __kernel void m13500_sxx (KERN_ATTR_ESALT (pstoken_t))
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx);
 
diff --git a/OpenCL/m13800_a1-pure.cl b/OpenCL/m13800_a1-pure.cl
index 7db54e87e..c148906b7 100644
--- a/OpenCL/m13800_a1-pure.cl
+++ b/OpenCL/m13800_a1-pure.cl
@@ -42,7 +42,7 @@ __kernel void m13800_mxx (KERN_ATTR_ESALT (win8phone_t))
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_update_global (&ctx, esalt_bufs[digests_offset].salt_buf, 128);
 
@@ -98,7 +98,7 @@ __kernel void m13800_sxx (KERN_ATTR_ESALT (win8phone_t))
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_update_global (&ctx, esalt_bufs[digests_offset].salt_buf, 128);
 
diff --git a/OpenCL/m13900_a1-pure.cl b/OpenCL/m13900_a1-pure.cl
index 791ec3d43..d9d9052d5 100644
--- a/OpenCL/m13900_a1-pure.cl
+++ b/OpenCL/m13900_a1-pure.cl
@@ -78,7 +78,7 @@ __kernel void m13900_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx1l;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
@@ -239,7 +239,7 @@ __kernel void m13900_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx1 = ctx1l;
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_final (&ctx1);
 
diff --git a/OpenCL/m14400_a1-pure.cl b/OpenCL/m14400_a1-pure.cl
index 10ba89ce4..3578a922f 100644
--- a/OpenCL/m14400_a1-pure.cl
+++ b/OpenCL/m14400_a1-pure.cl
@@ -141,7 +141,7 @@ __kernel void m14400_mxx (KERN_ATTR_BASIC ())
 
     sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len & 255);
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     d40[0] = 0x2d2d2d2d;
     d40[1] = 0;
@@ -231,7 +231,7 @@ __kernel void m14400_mxx (KERN_ATTR_BASIC ())
 
       sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
 
-      sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+      sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
       d40[0] = 0x2d2d2d2d;
       d40[1] = 0;
@@ -398,7 +398,7 @@ __kernel void m14400_sxx (KERN_ATTR_BASIC ())
 
     sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len & 255);
 
-    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     d40[0] = 0x2d2d2d2d;
     d40[1] = 0;
@@ -488,7 +488,7 @@ __kernel void m14400_sxx (KERN_ATTR_BASIC ())
 
       sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
 
-      sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+      sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
       d40[0] = 0x2d2d2d2d;
       d40[1] = 0;
diff --git a/OpenCL/m15000_a1-pure.cl b/OpenCL/m15000_a1-pure.cl
index 88e6e3111..899425593 100644
--- a/OpenCL/m15000_a1-pure.cl
+++ b/OpenCL/m15000_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m15000_mxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_update (&ctx, s, salt_len);
 
@@ -116,7 +116,7 @@ __kernel void m15000_sxx (KERN_ATTR_BASIC ())
   {
     sha512_ctx_t ctx = ctx0;
 
-    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha512_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m15500_a1-pure.cl b/OpenCL/m15500_a1-pure.cl
index da1b33afe..578806b63 100644
--- a/OpenCL/m15500_a1-pure.cl
+++ b/OpenCL/m15500_a1-pure.cl
@@ -51,7 +51,7 @@ __kernel void m15500_mxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_utf16be_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_utf16be_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx, s, salt_len);
 
@@ -122,7 +122,7 @@ __kernel void m15500_sxx (KERN_ATTR_BASIC ())
   {
     sha1_ctx_t ctx = ctx0;
 
-    sha1_update_global_utf16be_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha1_update_global_utf16be_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha1_update (&ctx, s, salt_len);
 
diff --git a/OpenCL/m16100_a1-pure.cl b/OpenCL/m16100_a1-pure.cl
index 68f7dadbc..31399deab 100644
--- a/OpenCL/m16100_a1-pure.cl
+++ b/OpenCL/m16100_a1-pure.cl
@@ -74,7 +74,7 @@ __kernel void m16100_mxx (KERN_ATTR_ESALT (tacacs_plus_t))
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     u32 sequence0[4];
     u32 sequence1[4];
@@ -232,7 +232,7 @@ __kernel void m16100_sxx (KERN_ATTR_ESALT (tacacs_plus_t))
   {
     md5_ctx_t ctx = ctx0;
 
-    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     u32 sequence0[4];
     u32 sequence1[4];
diff --git a/OpenCL/m16400_a1-pure.cl b/OpenCL/m16400_a1-pure.cl
index 366d77075..4100eed5f 100644
--- a/OpenCL/m16400_a1-pure.cl
+++ b/OpenCL/m16400_a1-pure.cl
@@ -204,7 +204,7 @@ __kernel void m16400_mxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    cram_md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    cram_md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     cram_md5_final (&ctx);
 
@@ -258,7 +258,7 @@ __kernel void m16400_sxx (KERN_ATTR_BASIC ())
   {
     md5_ctx_t ctx = ctx0;
 
-    cram_md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    cram_md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     cram_md5_final (&ctx);
 
diff --git a/OpenCL/m16600_a1-pure.cl b/OpenCL/m16600_a1-pure.cl
index 69bd4e2d5..6029c9aa3 100644
--- a/OpenCL/m16600_a1-pure.cl
+++ b/OpenCL/m16600_a1-pure.cl
@@ -95,7 +95,7 @@ __kernel void m16600_mxx (KERN_ATTR_ESALT (electrum_wallet_t))
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_final (&ctx);
 
@@ -268,7 +268,7 @@ __kernel void m16600_sxx (KERN_ATTR_ESALT (electrum_wallet_t))
   {
     sha256_ctx_t ctx = ctx0;
 
-    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     sha256_final (&ctx);
 
diff --git a/OpenCL/m18200_a1-pure.cl b/OpenCL/m18200_a1-pure.cl
index a1535dfe4..0c84a2980 100644
--- a/OpenCL/m18200_a1-pure.cl
+++ b/OpenCL/m18200_a1-pure.cl
@@ -413,7 +413,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_A
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);
 
@@ -473,7 +473,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_sxx (KERN_A
   {
     md4_ctx_t ctx = ctx0;
 
-    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
+    md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
 
     md4_final (&ctx);