From 0e428b3c407c8611dfa0d4b14a3bfaf8e23a0560 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Tue, 20 Nov 2018 15:32:41 +0100 Subject: [PATCH] Give the compiler a hint for automatic optimizations based on password length --- OpenCL/m00000_a1-pure.cl | 4 ++-- OpenCL/m00010_a1-pure.cl | 4 ++-- OpenCL/m00020_a1-pure.cl | 4 ++-- OpenCL/m00030_a1-pure.cl | 4 ++-- OpenCL/m00040_a1-pure.cl | 4 ++-- OpenCL/m00100_a1-pure.cl | 4 ++-- OpenCL/m00110_a1-pure.cl | 4 ++-- OpenCL/m00120_a1-pure.cl | 4 ++-- OpenCL/m00130_a1-pure.cl | 4 ++-- OpenCL/m00140_a1-pure.cl | 4 ++-- OpenCL/m00300_a1-pure.cl | 4 ++-- OpenCL/m00400-pure.cl | 2 +- OpenCL/m00900_a1-pure.cl | 4 ++-- OpenCL/m01000_a1-pure.cl | 4 ++-- OpenCL/m01100_a1-pure.cl | 4 ++-- OpenCL/m01300_a1-pure.cl | 4 ++-- OpenCL/m01400_a1-pure.cl | 4 ++-- OpenCL/m01410_a1-pure.cl | 4 ++-- OpenCL/m01420_a1-pure.cl | 4 ++-- OpenCL/m01430_a1-pure.cl | 4 ++-- OpenCL/m01440_a1-pure.cl | 4 ++-- OpenCL/m01700_a1-pure.cl | 4 ++-- OpenCL/m01710_a1-pure.cl | 4 ++-- OpenCL/m01720_a1-pure.cl | 4 ++-- OpenCL/m01730_a1-pure.cl | 4 ++-- OpenCL/m01740_a1-pure.cl | 4 ++-- OpenCL/m02100-pure.cl | 2 +- OpenCL/m02500-pure.cl | 2 +- OpenCL/m02610_a1-pure.cl | 4 ++-- OpenCL/m02810_a1-pure.cl | 4 ++-- OpenCL/m03710_a1-pure.cl | 4 ++-- OpenCL/m03800_a1-pure.cl | 4 ++-- OpenCL/m03910_a1-pure.cl | 4 ++-- OpenCL/m04010_a1-pure.cl | 4 ++-- OpenCL/m04110_a1-pure.cl | 4 ++-- OpenCL/m04310_a1-pure.cl | 4 ++-- OpenCL/m04400_a1-pure.cl | 4 ++-- OpenCL/m04500_a1-pure.cl | 4 ++-- OpenCL/m04520_a1-pure.cl | 4 ++-- OpenCL/m04700_a1-pure.cl | 4 ++-- OpenCL/m04800_a1-pure.cl | 4 ++-- OpenCL/m04900_a1-pure.cl | 4 ++-- OpenCL/m05100_a1-pure.cl | 4 ++-- OpenCL/m05200-pure.cl | 2 +- OpenCL/m05500_a1-pure.cl | 4 ++-- OpenCL/m05600_a1-pure.cl | 4 ++-- OpenCL/m05800-pure.cl | 2 +- OpenCL/m06000_a1-pure.cl | 4 ++-- OpenCL/m06100_a1-pure.cl | 4 ++-- OpenCL/m06400-pure.cl | 2 +- OpenCL/m06500-pure.cl | 2 +- OpenCL/m06600-pure.cl | 2 +- OpenCL/m06700-pure.cl | 2 +- OpenCL/m06800-pure.cl | 2 +- OpenCL/m07000_a1-pure.cl | 4 ++-- OpenCL/m07100-pure.cl | 2 +- OpenCL/m07500_a1-pure.cl | 4 ++-- OpenCL/m07900-pure.cl | 2 +- OpenCL/m08100_a1-pure.cl | 4 ++-- OpenCL/m08200-pure.cl | 2 +- OpenCL/m08300_a1-pure.cl | 4 ++-- OpenCL/m08400_a1-pure.cl | 4 ++-- OpenCL/m08800-pure.cl | 2 +- OpenCL/m08900-pure.cl | 4 ++-- OpenCL/m09400-pure.cl | 2 +- OpenCL/m09500-pure.cl | 2 +- OpenCL/m09600-pure.cl | 2 +- OpenCL/m09900_a1-pure.cl | 4 ++-- OpenCL/m10300-pure.cl | 4 ++-- OpenCL/m10700-optimized.cl | 2 +- OpenCL/m10700-pure.cl | 2 +- OpenCL/m10800_a1-pure.cl | 4 ++-- OpenCL/m10900-pure.cl | 2 +- OpenCL/m11000_a1-pure.cl | 4 ++-- OpenCL/m11100_a1-pure.cl | 4 ++-- OpenCL/m11200_a1-pure.cl | 4 ++-- OpenCL/m11300-pure.cl | 2 +- OpenCL/m11400_a1-pure.cl | 4 ++-- OpenCL/m11700_a1-pure.cl | 4 ++-- OpenCL/m11800_a1-pure.cl | 4 ++-- OpenCL/m11900-pure.cl | 2 +- OpenCL/m12000-pure.cl | 2 +- OpenCL/m12200-pure.cl | 2 +- OpenCL/m12300-pure.cl | 2 +- OpenCL/m12600_a1-pure.cl | 4 ++-- OpenCL/m12700-pure.cl | 2 +- OpenCL/m12800-pure.cl | 2 +- OpenCL/m12900-pure.cl | 2 +- OpenCL/m13000-pure.cl | 2 +- OpenCL/m13100_a1-pure.cl | 4 ++-- OpenCL/m13200-pure.cl | 2 +- OpenCL/m13300_a1-pure.cl | 4 ++-- OpenCL/m13400-pure.cl | 2 +- OpenCL/m13500_a1-pure.cl | 4 ++-- OpenCL/m13600-pure.cl | 2 +- OpenCL/m13800_a1-pure.cl | 4 ++-- OpenCL/m13900_a1-pure.cl | 4 ++-- OpenCL/m14400_a1-pure.cl | 8 ++++---- OpenCL/m14611-pure.cl | 2 +- OpenCL/m14612-pure.cl | 2 +- OpenCL/m14613-pure.cl | 2 +- OpenCL/m14621-pure.cl | 2 +- OpenCL/m14622-pure.cl | 2 +- OpenCL/m14623-pure.cl | 2 +- OpenCL/m14631-pure.cl | 2 +- OpenCL/m14632-pure.cl | 2 +- OpenCL/m14633-pure.cl | 2 +- OpenCL/m14641-pure.cl | 2 +- OpenCL/m14642-pure.cl | 2 +- OpenCL/m14643-pure.cl | 2 +- OpenCL/m14700-pure.cl | 2 +- OpenCL/m14800-pure.cl | 2 +- OpenCL/m15000_a1-pure.cl | 4 ++-- OpenCL/m15100-pure.cl | 2 +- OpenCL/m15300-pure.cl | 4 ++-- OpenCL/m15500_a1-pure.cl | 4 ++-- OpenCL/m15600-pure.cl | 2 +- OpenCL/m15700-pure.cl | 4 ++-- OpenCL/m15900-pure.cl | 4 ++-- OpenCL/m16100_a1-pure.cl | 4 ++-- OpenCL/m16200-pure.cl | 2 +- OpenCL/m16300-pure.cl | 4 ++-- OpenCL/m16400_a1-pure.cl | 4 ++-- OpenCL/m16600_a1-pure.cl | 4 ++-- OpenCL/m16800-pure.cl | 2 +- OpenCL/m16900-pure.cl | 2 +- OpenCL/m18200_a1-pure.cl | 4 ++-- OpenCL/m18300-pure.cl | 2 +- 128 files changed, 206 insertions(+), 206 deletions(-) diff --git a/OpenCL/m00000_a1-pure.cl b/OpenCL/m00000_a1-pure.cl index 3a9c95472..fef760853 100644 --- a/OpenCL/m00000_a1-pure.cl +++ b/OpenCL/m00000_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m00000_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m00000_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00010_a1-pure.cl b/OpenCL/m00010_a1-pure.cl index 40373293f..8ae461f10 100644 --- a/OpenCL/m00010_a1-pure.cl +++ b/OpenCL/m00010_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m00010_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m00010_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00020_a1-pure.cl b/OpenCL/m00020_a1-pure.cl index 540224ad5..e862c48eb 100644 --- a/OpenCL/m00020_a1-pure.cl +++ b/OpenCL/m00020_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m00020_mxx (KERN_ATTR_BASIC ()) md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m00020_sxx (KERN_ATTR_BASIC ()) md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00030_a1-pure.cl b/OpenCL/m00030_a1-pure.cl index f4e2fd4a1..83ed6766a 100644 --- a/OpenCL/m00030_a1-pure.cl +++ b/OpenCL/m00030_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m00030_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m00030_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00040_a1-pure.cl b/OpenCL/m00040_a1-pure.cl index b1424dcdd..67fa78329 100644 --- a/OpenCL/m00040_a1-pure.cl +++ b/OpenCL/m00040_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m00040_mxx (KERN_ATTR_BASIC ()) md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m00040_sxx (KERN_ATTR_BASIC ()) md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00100_a1-pure.cl b/OpenCL/m00100_a1-pure.cl index f7ff108f6..375bedbd7 100644 --- a/OpenCL/m00100_a1-pure.cl +++ b/OpenCL/m00100_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m00100_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m00100_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00110_a1-pure.cl b/OpenCL/m00110_a1-pure.cl index ad8bac633..2bc8177ca 100644 --- a/OpenCL/m00110_a1-pure.cl +++ b/OpenCL/m00110_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m00110_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m00110_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00120_a1-pure.cl b/OpenCL/m00120_a1-pure.cl index 5b574c1cf..8acbd83b4 100644 --- a/OpenCL/m00120_a1-pure.cl +++ b/OpenCL/m00120_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m00120_mxx (KERN_ATTR_BASIC ()) sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m00120_sxx (KERN_ATTR_BASIC ()) sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00130_a1-pure.cl b/OpenCL/m00130_a1-pure.cl index 5175b3d99..a9b17120c 100644 --- a/OpenCL/m00130_a1-pure.cl +++ b/OpenCL/m00130_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m00130_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m00130_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00140_a1-pure.cl b/OpenCL/m00140_a1-pure.cl index ed211f9ae..02842ffc0 100644 --- a/OpenCL/m00140_a1-pure.cl +++ b/OpenCL/m00140_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m00140_mxx (KERN_ATTR_BASIC ()) sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m00140_sxx (KERN_ATTR_BASIC ()) sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00300_a1-pure.cl b/OpenCL/m00300_a1-pure.cl index 57a25e292..67edd34b8 100644 --- a/OpenCL/m00300_a1-pure.cl +++ b/OpenCL/m00300_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m00300_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -109,7 +109,7 @@ __kernel void m00300_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m00400-pure.cl b/OpenCL/m00400-pure.cl index 85515b503..2f215a53b 100644 --- a/OpenCL/m00400-pure.cl +++ b/OpenCL/m00400-pure.cl @@ -36,7 +36,7 @@ __kernel void m00400_init (KERN_ATTR_TMPS (phpass_tmp_t)) md5_update_global (&md5_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len); + md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len & 255); md5_final (&md5_ctx); diff --git a/OpenCL/m00900_a1-pure.cl b/OpenCL/m00900_a1-pure.cl index 996a8937b..50118c203 100644 --- a/OpenCL/m00900_a1-pure.cl +++ b/OpenCL/m00900_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m00900_mxx (KERN_ATTR_BASIC ()) md4_init (&ctx0); - md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m00900_sxx (KERN_ATTR_BASIC ()) md4_init (&ctx0); - md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01000_a1-pure.cl b/OpenCL/m01000_a1-pure.cl index 67b362099..c5b438e0b 100644 --- a/OpenCL/m01000_a1-pure.cl +++ b/OpenCL/m01000_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m01000_mxx (KERN_ATTR_BASIC ()) md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m01000_sxx (KERN_ATTR_BASIC ()) md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01100_a1-pure.cl b/OpenCL/m01100_a1-pure.cl index 9167ef725..a6f8b7c81 100644 --- a/OpenCL/m01100_a1-pure.cl +++ b/OpenCL/m01100_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m01100_mxx (KERN_ATTR_BASIC ()) md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -119,7 +119,7 @@ __kernel void m01100_sxx (KERN_ATTR_BASIC ()) md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01300_a1-pure.cl b/OpenCL/m01300_a1-pure.cl index 5220a9de1..fcb155729 100644 --- a/OpenCL/m01300_a1-pure.cl +++ b/OpenCL/m01300_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m01300_mxx (KERN_ATTR_BASIC ()) sha224_init (&ctx0); - sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m01300_sxx (KERN_ATTR_BASIC ()) sha224_init (&ctx0); - sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01400_a1-pure.cl b/OpenCL/m01400_a1-pure.cl index 2aea28843..fc040002b 100644 --- a/OpenCL/m01400_a1-pure.cl +++ b/OpenCL/m01400_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m01400_mxx (KERN_ATTR_BASIC ()) sha256_init (&ctx0); - sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m01400_sxx (KERN_ATTR_BASIC ()) sha256_init (&ctx0); - sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01410_a1-pure.cl b/OpenCL/m01410_a1-pure.cl index c0cf80209..eaa6cbae0 100644 --- a/OpenCL/m01410_a1-pure.cl +++ b/OpenCL/m01410_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m01410_mxx (KERN_ATTR_BASIC ()) sha256_init (&ctx0); - sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m01410_sxx (KERN_ATTR_BASIC ()) sha256_init (&ctx0); - sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01420_a1-pure.cl b/OpenCL/m01420_a1-pure.cl index b5e89184a..429e62e2e 100644 --- a/OpenCL/m01420_a1-pure.cl +++ b/OpenCL/m01420_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m01420_mxx (KERN_ATTR_BASIC ()) sha256_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m01420_sxx (KERN_ATTR_BASIC ()) sha256_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01430_a1-pure.cl b/OpenCL/m01430_a1-pure.cl index 7f804c0d0..1f441e83c 100644 --- a/OpenCL/m01430_a1-pure.cl +++ b/OpenCL/m01430_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m01430_mxx (KERN_ATTR_BASIC ()) sha256_init (&ctx0); - sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m01430_sxx (KERN_ATTR_BASIC ()) sha256_init (&ctx0); - sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01440_a1-pure.cl b/OpenCL/m01440_a1-pure.cl index 461b59443..c3e7070b4 100644 --- a/OpenCL/m01440_a1-pure.cl +++ b/OpenCL/m01440_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m01440_mxx (KERN_ATTR_BASIC ()) sha256_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m01440_sxx (KERN_ATTR_BASIC ()) sha256_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01700_a1-pure.cl b/OpenCL/m01700_a1-pure.cl index b066bde73..038558998 100644 --- a/OpenCL/m01700_a1-pure.cl +++ b/OpenCL/m01700_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m01700_mxx (KERN_ATTR_BASIC ()) sha512_init (&ctx0); - sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m01700_sxx (KERN_ATTR_BASIC ()) sha512_init (&ctx0); - sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01710_a1-pure.cl b/OpenCL/m01710_a1-pure.cl index e2c3572aa..296a74027 100644 --- a/OpenCL/m01710_a1-pure.cl +++ b/OpenCL/m01710_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m01710_mxx (KERN_ATTR_BASIC ()) sha512_init (&ctx0); - sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m01710_sxx (KERN_ATTR_BASIC ()) sha512_init (&ctx0); - sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01720_a1-pure.cl b/OpenCL/m01720_a1-pure.cl index 31e8be3f2..3fa6901f4 100644 --- a/OpenCL/m01720_a1-pure.cl +++ b/OpenCL/m01720_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m01720_mxx (KERN_ATTR_BASIC ()) sha512_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m01720_sxx (KERN_ATTR_BASIC ()) sha512_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01730_a1-pure.cl b/OpenCL/m01730_a1-pure.cl index 2e0b25f8a..77c856527 100644 --- a/OpenCL/m01730_a1-pure.cl +++ b/OpenCL/m01730_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m01730_mxx (KERN_ATTR_BASIC ()) sha512_init (&ctx0); - sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m01730_sxx (KERN_ATTR_BASIC ()) sha512_init (&ctx0); - sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m01740_a1-pure.cl b/OpenCL/m01740_a1-pure.cl index 64466a77e..0f0d9692b 100644 --- a/OpenCL/m01740_a1-pure.cl +++ b/OpenCL/m01740_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m01740_mxx (KERN_ATTR_BASIC ()) sha512_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m01740_sxx (KERN_ATTR_BASIC ()) sha512_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m02100-pure.cl b/OpenCL/m02100-pure.cl index ecaa9b75b..08db170e9 100644 --- a/OpenCL/m02100-pure.cl +++ b/OpenCL/m02100-pure.cl @@ -69,7 +69,7 @@ __kernel void m02100_init (KERN_ATTR_TMPS (dcc2_tmp_t)) md4_init (&md4_ctx1); - md4_update_global_utf16le (&md4_ctx1, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&md4_ctx1, pws[gid].i, pws[gid].pw_len & 255); md4_final (&md4_ctx1); diff --git a/OpenCL/m02500-pure.cl b/OpenCL/m02500-pure.cl index 1e1ef1691..5d66edf03 100644 --- a/OpenCL/m02500-pure.cl +++ b/OpenCL/m02500-pure.cl @@ -97,7 +97,7 @@ __kernel void m02500_init (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m02610_a1-pure.cl b/OpenCL/m02610_a1-pure.cl index 9fa3f4ee4..a36061bf5 100644 --- a/OpenCL/m02610_a1-pure.cl +++ b/OpenCL/m02610_a1-pure.cl @@ -71,7 +71,7 @@ __kernel void m02610_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -184,7 +184,7 @@ __kernel void m02610_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m02810_a1-pure.cl b/OpenCL/m02810_a1-pure.cl index 47653774b..17dc88fbe 100644 --- a/OpenCL/m02810_a1-pure.cl +++ b/OpenCL/m02810_a1-pure.cl @@ -71,7 +71,7 @@ __kernel void m02810_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -207,7 +207,7 @@ __kernel void m02810_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m03710_a1-pure.cl b/OpenCL/m03710_a1-pure.cl index e6ed23e45..5455decfa 100644 --- a/OpenCL/m03710_a1-pure.cl +++ b/OpenCL/m03710_a1-pure.cl @@ -71,7 +71,7 @@ __kernel void m03710_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -197,7 +197,7 @@ __kernel void m03710_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m03800_a1-pure.cl b/OpenCL/m03800_a1-pure.cl index 01197f565..fecfd6556 100644 --- a/OpenCL/m03800_a1-pure.cl +++ b/OpenCL/m03800_a1-pure.cl @@ -43,7 +43,7 @@ __kernel void m03800_mxx (KERN_ATTR_BASIC ()) md5_update (&ctx0, s, salt_len); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -110,7 +110,7 @@ __kernel void m03800_sxx (KERN_ATTR_BASIC ()) md5_update (&ctx0, s, salt_len); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m03910_a1-pure.cl b/OpenCL/m03910_a1-pure.cl index c74edca4f..b5f36d7a2 100644 --- a/OpenCL/m03910_a1-pure.cl +++ b/OpenCL/m03910_a1-pure.cl @@ -71,7 +71,7 @@ __kernel void m03910_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -207,7 +207,7 @@ __kernel void m03910_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04010_a1-pure.cl b/OpenCL/m04010_a1-pure.cl index cc980eafe..7a8998795 100644 --- a/OpenCL/m04010_a1-pure.cl +++ b/OpenCL/m04010_a1-pure.cl @@ -66,7 +66,7 @@ __kernel void m04010_mxx (KERN_ATTR_BASIC ()) md5_ctx_t ctx0t = ctx0; - md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -183,7 +183,7 @@ __kernel void m04010_sxx (KERN_ATTR_BASIC ()) md5_ctx_t ctx0t = ctx0; - md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04110_a1-pure.cl b/OpenCL/m04110_a1-pure.cl index c22847369..1f3c9af3a 100644 --- a/OpenCL/m04110_a1-pure.cl +++ b/OpenCL/m04110_a1-pure.cl @@ -77,7 +77,7 @@ __kernel void m04110_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0t); - md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -207,7 +207,7 @@ __kernel void m04110_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0t); - md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04310_a1-pure.cl b/OpenCL/m04310_a1-pure.cl index dfa730486..37ba88f15 100644 --- a/OpenCL/m04310_a1-pure.cl +++ b/OpenCL/m04310_a1-pure.cl @@ -71,7 +71,7 @@ __kernel void m04310_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -184,7 +184,7 @@ __kernel void m04310_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04400_a1-pure.cl b/OpenCL/m04400_a1-pure.cl index 25d4e9f6e..4841e5398 100644 --- a/OpenCL/m04400_a1-pure.cl +++ b/OpenCL/m04400_a1-pure.cl @@ -63,7 +63,7 @@ __kernel void m04400_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -170,7 +170,7 @@ __kernel void m04400_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04500_a1-pure.cl b/OpenCL/m04500_a1-pure.cl index 9fd6c48a3..fe9ce3ce3 100644 --- a/OpenCL/m04500_a1-pure.cl +++ b/OpenCL/m04500_a1-pure.cl @@ -62,7 +62,7 @@ __kernel void m04500_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -169,7 +169,7 @@ __kernel void m04500_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04520_a1-pure.cl b/OpenCL/m04520_a1-pure.cl index 3424ac2bb..b5488cbca 100644 --- a/OpenCL/m04520_a1-pure.cl +++ b/OpenCL/m04520_a1-pure.cl @@ -68,7 +68,7 @@ __kernel void m04520_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx1l); - sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -190,7 +190,7 @@ __kernel void m04520_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx1l); - sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04700_a1-pure.cl b/OpenCL/m04700_a1-pure.cl index 09208e505..7da0e4f5f 100644 --- a/OpenCL/m04700_a1-pure.cl +++ b/OpenCL/m04700_a1-pure.cl @@ -63,7 +63,7 @@ __kernel void m04700_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -165,7 +165,7 @@ __kernel void m04700_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04800_a1-pure.cl b/OpenCL/m04800_a1-pure.cl index 69bc376b1..6e527c17f 100644 --- a/OpenCL/m04800_a1-pure.cl +++ b/OpenCL/m04800_a1-pure.cl @@ -45,7 +45,7 @@ __kernel void m04800_mxx (KERN_ATTR_BASIC ()) ctx0.len = 1; - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -114,7 +114,7 @@ __kernel void m04800_sxx (KERN_ATTR_BASIC ()) ctx0.len = 1; - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m04900_a1-pure.cl b/OpenCL/m04900_a1-pure.cl index c09d76df9..17c62a3ea 100644 --- a/OpenCL/m04900_a1-pure.cl +++ b/OpenCL/m04900_a1-pure.cl @@ -43,7 +43,7 @@ __kernel void m04900_mxx (KERN_ATTR_BASIC ()) sha1_update (&ctx0, s, salt_len); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -110,7 +110,7 @@ __kernel void m04900_sxx (KERN_ATTR_BASIC ()) sha1_update (&ctx0, s, salt_len); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m05100_a1-pure.cl b/OpenCL/m05100_a1-pure.cl index 0cffa93ca..ceac4ee19 100644 --- a/OpenCL/m05100_a1-pure.cl +++ b/OpenCL/m05100_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m05100_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m05100_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m05200-pure.cl b/OpenCL/m05200-pure.cl index c659110cf..6b37b9a72 100644 --- a/OpenCL/m05200-pure.cl +++ b/OpenCL/m05200-pure.cl @@ -30,7 +30,7 @@ __kernel void m05200_init (KERN_ATTR_TMPS (pwsafe3_tmp_t)) sha256_init (&ctx); - sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha256_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); diff --git a/OpenCL/m05500_a1-pure.cl b/OpenCL/m05500_a1-pure.cl index da550d0b3..a09d65d30 100644 --- a/OpenCL/m05500_a1-pure.cl +++ b/OpenCL/m05500_a1-pure.cl @@ -545,7 +545,7 @@ __kernel void m05500_mxx (KERN_ATTR_BASIC ()) md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -679,7 +679,7 @@ __kernel void m05500_sxx (KERN_ATTR_BASIC ()) md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m05600_a1-pure.cl b/OpenCL/m05600_a1-pure.cl index 2f1204932..581389943 100644 --- a/OpenCL/m05600_a1-pure.cl +++ b/OpenCL/m05600_a1-pure.cl @@ -33,7 +33,7 @@ __kernel void m05600_mxx (KERN_ATTR_ESALT (netntlm_t)) md4_init (&ctx10); - md4_update_global_utf16le (&ctx10, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx10, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -142,7 +142,7 @@ __kernel void m05600_sxx (KERN_ATTR_ESALT (netntlm_t)) md4_init (&ctx10); - md4_update_global_utf16le (&ctx10, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx10, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m05800-pure.cl b/OpenCL/m05800-pure.cl index 434f6bc3e..f2ca70444 100644 --- a/OpenCL/m05800-pure.cl +++ b/OpenCL/m05800-pure.cl @@ -2214,7 +2214,7 @@ __kernel void m05800_init (KERN_ATTR_TMPS (androidpin_tmp_t)) sha1_update_64 (&ctx, w0, w1, w2, w3, 1); - sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha1_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); diff --git a/OpenCL/m06000_a1-pure.cl b/OpenCL/m06000_a1-pure.cl index 37c33282d..4fa59fe25 100644 --- a/OpenCL/m06000_a1-pure.cl +++ b/OpenCL/m06000_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m06000_mxx (KERN_ATTR_BASIC ()) ripemd160_init (&ctx0); - ripemd160_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + ripemd160_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m06000_sxx (KERN_ATTR_BASIC ()) ripemd160_init (&ctx0); - ripemd160_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + ripemd160_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m06100_a1-pure.cl b/OpenCL/m06100_a1-pure.cl index 66e5b6208..16cbec86b 100644 --- a/OpenCL/m06100_a1-pure.cl +++ b/OpenCL/m06100_a1-pure.cl @@ -63,7 +63,7 @@ __kernel void m06100_mxx (KERN_ATTR_BASIC ()) whirlpool_init (&ctx0, s_Ch, s_Cl); - whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -148,7 +148,7 @@ __kernel void m06100_sxx (KERN_ATTR_BASIC ()) whirlpool_init (&ctx0, s_Ch, s_Cl); - whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m06400-pure.cl b/OpenCL/m06400-pure.cl index a312bf990..8baf1652c 100644 --- a/OpenCL/m06400-pure.cl +++ b/OpenCL/m06400-pure.cl @@ -70,7 +70,7 @@ __kernel void m06400_init (KERN_ATTR_TMPS (sha256aix_tmp_t)) sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m06500-pure.cl b/OpenCL/m06500-pure.cl index acb70f006..97e116ab0 100644 --- a/OpenCL/m06500-pure.cl +++ b/OpenCL/m06500-pure.cl @@ -86,7 +86,7 @@ __kernel void m06500_init (KERN_ATTR_TMPS (sha512aix_tmp_t)) sha512_hmac_ctx_t sha512_hmac_ctx; - sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m06600-pure.cl b/OpenCL/m06600-pure.cl index badea394f..a13cc962d 100644 --- a/OpenCL/m06600-pure.cl +++ b/OpenCL/m06600-pure.cl @@ -65,7 +65,7 @@ __kernel void m06600_init (KERN_ATTR_TMPS (agilekey_tmp_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m06700-pure.cl b/OpenCL/m06700-pure.cl index 7a6fb736f..5afd6c37b 100644 --- a/OpenCL/m06700-pure.cl +++ b/OpenCL/m06700-pure.cl @@ -64,7 +64,7 @@ __kernel void m06700_init (KERN_ATTR_TMPS (sha1aix_tmp_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m06800-pure.cl b/OpenCL/m06800-pure.cl index 0bc104740..c20844969 100644 --- a/OpenCL/m06800-pure.cl +++ b/OpenCL/m06800-pure.cl @@ -71,7 +71,7 @@ __kernel void m06800_init (KERN_ATTR_TMPS (lastpass_tmp_t)) sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m07000_a1-pure.cl b/OpenCL/m07000_a1-pure.cl index 66d9a1325..521b13d63 100644 --- a/OpenCL/m07000_a1-pure.cl +++ b/OpenCL/m07000_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m07000_mxx (KERN_ATTR_BASIC ()) sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -118,7 +118,7 @@ __kernel void m07000_sxx (KERN_ATTR_BASIC ()) sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m07100-pure.cl b/OpenCL/m07100-pure.cl index 3651b1dd0..a01f0a03a 100644 --- a/OpenCL/m07100-pure.cl +++ b/OpenCL/m07100-pure.cl @@ -86,7 +86,7 @@ __kernel void m07100_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, pbkdf2_sha sha512_hmac_ctx_t sha512_hmac_ctx; - sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m07500_a1-pure.cl b/OpenCL/m07500_a1-pure.cl index 7109a9a59..c80345831 100644 --- a/OpenCL/m07500_a1-pure.cl +++ b/OpenCL/m07500_a1-pure.cl @@ -307,7 +307,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_A md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -376,7 +376,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_sxx (KERN_A md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m07900-pure.cl b/OpenCL/m07900-pure.cl index f33506e08..b576fbbf9 100644 --- a/OpenCL/m07900-pure.cl +++ b/OpenCL/m07900-pure.cl @@ -29,7 +29,7 @@ __kernel void m07900_init (KERN_ATTR_TMPS (drupal7_tmp_t)) sha512_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha512_final (&ctx); diff --git a/OpenCL/m08100_a1-pure.cl b/OpenCL/m08100_a1-pure.cl index 395b55c8f..20cb6d5e6 100644 --- a/OpenCL/m08100_a1-pure.cl +++ b/OpenCL/m08100_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m08100_mxx (KERN_ATTR_BASIC ()) sha1_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m08100_sxx (KERN_ATTR_BASIC ()) sha1_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m08200-pure.cl b/OpenCL/m08200-pure.cl index 5be1abfc0..4eeb2d753 100644 --- a/OpenCL/m08200-pure.cl +++ b/OpenCL/m08200-pure.cl @@ -87,7 +87,7 @@ __kernel void m08200_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, cloudkey_t sha512_hmac_ctx_t sha512_hmac_ctx; - sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m08300_a1-pure.cl b/OpenCL/m08300_a1-pure.cl index 05390dc3d..27a24da49 100644 --- a/OpenCL/m08300_a1-pure.cl +++ b/OpenCL/m08300_a1-pure.cl @@ -62,7 +62,7 @@ __kernel void m08300_mxx (KERN_ATTR_BASIC ()) ctx1.len = 1; - sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len); + 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); @@ -177,7 +177,7 @@ __kernel void m08300_sxx (KERN_ATTR_BASIC ()) ctx1.len = 1; - sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len); + 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); diff --git a/OpenCL/m08400_a1-pure.cl b/OpenCL/m08400_a1-pure.cl index e3b7c674e..6a80b86dd 100644 --- a/OpenCL/m08400_a1-pure.cl +++ b/OpenCL/m08400_a1-pure.cl @@ -68,7 +68,7 @@ __kernel void m08400_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx1l); - sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -229,7 +229,7 @@ __kernel void m08400_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx1l); - sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m08800-pure.cl b/OpenCL/m08800-pure.cl index 8f41e9f23..08762bdc4 100644 --- a/OpenCL/m08800-pure.cl +++ b/OpenCL/m08800-pure.cl @@ -63,7 +63,7 @@ __kernel void m08800_init (KERN_ATTR_TMPS_ESALT (androidfde_tmp_t, androidfde_t) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m08900-pure.cl b/OpenCL/m08900-pure.cl index fd388c70d..1b6cc9d32 100644 --- a/OpenCL/m08900-pure.cl +++ b/OpenCL/m08900-pure.cl @@ -219,7 +219,7 @@ __kernel void m08900_init (KERN_ATTR_TMPS (scrypt_tmp_t)) sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); @@ -327,7 +327,7 @@ __kernel void m08900_comp (KERN_ATTR_TMPS (scrypt_tmp_t)) sha256_hmac_ctx_t ctx; - sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); for (u32 l = 0; l < SCRYPT_CNT4; l += 4) { diff --git a/OpenCL/m09400-pure.cl b/OpenCL/m09400-pure.cl index c8902b85c..78773a829 100644 --- a/OpenCL/m09400-pure.cl +++ b/OpenCL/m09400-pure.cl @@ -33,7 +33,7 @@ __kernel void m09400_init (KERN_ATTR_TMPS_ESALT (office2007_tmp_t, office2007_t) sha1_update_global (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha1_final (&ctx); diff --git a/OpenCL/m09500-pure.cl b/OpenCL/m09500-pure.cl index 068f496a0..a0d3e4025 100644 --- a/OpenCL/m09500-pure.cl +++ b/OpenCL/m09500-pure.cl @@ -33,7 +33,7 @@ __kernel void m09500_init (KERN_ATTR_TMPS_ESALT (office2010_tmp_t, office2010_t) sha1_update_global (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha1_final (&ctx); diff --git a/OpenCL/m09600-pure.cl b/OpenCL/m09600-pure.cl index a77319f43..4f59fa484 100644 --- a/OpenCL/m09600-pure.cl +++ b/OpenCL/m09600-pure.cl @@ -33,7 +33,7 @@ __kernel void m09600_init (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t) sha512_update_global (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha512_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha512_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha512_final (&ctx); diff --git a/OpenCL/m09900_a1-pure.cl b/OpenCL/m09900_a1-pure.cl index faa92d443..d22ce547e 100644 --- a/OpenCL/m09900_a1-pure.cl +++ b/OpenCL/m09900_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m09900_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m09900_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m10300-pure.cl b/OpenCL/m10300-pure.cl index 6541ce2b2..0418df911 100644 --- a/OpenCL/m10300-pure.cl +++ b/OpenCL/m10300-pure.cl @@ -27,7 +27,7 @@ __kernel void m10300_init (KERN_ATTR_TMPS (saph_sha1_tmp_t)) sha1_init (&ctx); - sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha1_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); @@ -58,7 +58,7 @@ __kernel void m10300_loop (KERN_ATTR_TMPS (saph_sha1_tmp_t)) sha1_init (&ctx); - sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); /** * load diff --git a/OpenCL/m10700-optimized.cl b/OpenCL/m10700-optimized.cl index 491fd0dee..87aa1df61 100644 --- a/OpenCL/m10700-optimized.cl +++ b/OpenCL/m10700-optimized.cl @@ -529,7 +529,7 @@ __kernel void m10700_init (KERN_ATTR_TMPS_ESALT (pdf17l8_tmp_t, pdf_t)) sha256_init (&ctx); - sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha256_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); diff --git a/OpenCL/m10700-pure.cl b/OpenCL/m10700-pure.cl index 373faed69..04c595c22 100644 --- a/OpenCL/m10700-pure.cl +++ b/OpenCL/m10700-pure.cl @@ -1118,7 +1118,7 @@ __kernel void m10700_init (KERN_ATTR_TMPS_ESALT (pdf17l8_tmp_t, pdf_t)) sha256_init (&ctx); - sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha256_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); diff --git a/OpenCL/m10800_a1-pure.cl b/OpenCL/m10800_a1-pure.cl index 84e1655bc..267ffac25 100644 --- a/OpenCL/m10800_a1-pure.cl +++ b/OpenCL/m10800_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m10800_mxx (KERN_ATTR_BASIC ()) sha384_init (&ctx0); - sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -86,7 +86,7 @@ __kernel void m10800_sxx (KERN_ATTR_BASIC ()) sha384_init (&ctx0); - sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m10900-pure.cl b/OpenCL/m10900-pure.cl index c228d3403..2d88235a2 100644 --- a/OpenCL/m10900-pure.cl +++ b/OpenCL/m10900-pure.cl @@ -70,7 +70,7 @@ __kernel void m10900_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m11000_a1-pure.cl b/OpenCL/m11000_a1-pure.cl index 09a57f290..38a9cce36 100644 --- a/OpenCL/m11000_a1-pure.cl +++ b/OpenCL/m11000_a1-pure.cl @@ -34,7 +34,7 @@ __kernel void m11000_mxx (KERN_ATTR_BASIC ()) md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -90,7 +90,7 @@ __kernel void m11000_sxx (KERN_ATTR_BASIC ()) md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m11100_a1-pure.cl b/OpenCL/m11100_a1-pure.cl index 8884905e8..6046dd6e7 100644 --- a/OpenCL/m11100_a1-pure.cl +++ b/OpenCL/m11100_a1-pure.cl @@ -84,7 +84,7 @@ __kernel void m11100_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0t); - md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -240,7 +240,7 @@ __kernel void m11100_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0t); - md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m11200_a1-pure.cl b/OpenCL/m11200_a1-pure.cl index 10426c19d..eedfe51ea 100644 --- a/OpenCL/m11200_a1-pure.cl +++ b/OpenCL/m11200_a1-pure.cl @@ -38,7 +38,7 @@ __kernel void m11200_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx2l); - sha1_update_global_swap (&ctx2l, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx2l, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -164,7 +164,7 @@ __kernel void m11200_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx2l); - sha1_update_global_swap (&ctx2l, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx2l, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m11300-pure.cl b/OpenCL/m11300-pure.cl index 54ac4c55f..993d2b724 100644 --- a/OpenCL/m11300-pure.cl +++ b/OpenCL/m11300-pure.cl @@ -86,7 +86,7 @@ __kernel void m11300_init (KERN_ATTR_TMPS_ESALT (bitcoin_wallet_tmp_t, bitcoin_w sha512_init (&ctx); - sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha512_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); diff --git a/OpenCL/m11400_a1-pure.cl b/OpenCL/m11400_a1-pure.cl index 0643a3eb8..7067a3325 100644 --- a/OpenCL/m11400_a1-pure.cl +++ b/OpenCL/m11400_a1-pure.cl @@ -64,7 +64,7 @@ __kernel void m11400_mxx (KERN_ATTR_ESALT (sip_t)) md5_update_global (&ctx0, esalt_bufs[digests_offset].salt_buf, esalt_bufs[digests_offset].salt_len); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -170,7 +170,7 @@ __kernel void m11400_sxx (KERN_ATTR_ESALT (sip_t)) md5_update_global (&ctx0, esalt_bufs[digests_offset].salt_buf, esalt_bufs[digests_offset].salt_len); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m11700_a1-pure.cl b/OpenCL/m11700_a1-pure.cl index e6d531ff8..8b05a6969 100644 --- a/OpenCL/m11700_a1-pure.cl +++ b/OpenCL/m11700_a1-pure.cl @@ -61,7 +61,7 @@ __kernel void m11700_mxx (KERN_ATTR_BASIC ()) streebog256_init (&ctx0, s_sbob_sl64); - streebog256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + streebog256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -144,7 +144,7 @@ __kernel void m11700_sxx (KERN_ATTR_BASIC ()) streebog256_init (&ctx0, s_sbob_sl64); - streebog256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + streebog256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m11800_a1-pure.cl b/OpenCL/m11800_a1-pure.cl index 770c0ea7b..e4c63e3f9 100644 --- a/OpenCL/m11800_a1-pure.cl +++ b/OpenCL/m11800_a1-pure.cl @@ -61,7 +61,7 @@ __kernel void m11800_mxx (KERN_ATTR_BASIC ()) streebog512_init (&ctx0, s_sbob_sl64); - streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -144,7 +144,7 @@ __kernel void m11800_sxx (KERN_ATTR_BASIC ()) streebog512_init (&ctx0, s_sbob_sl64); - streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m11900-pure.cl b/OpenCL/m11900-pure.cl index 12e00ed2f..97db81c87 100644 --- a/OpenCL/m11900-pure.cl +++ b/OpenCL/m11900-pure.cl @@ -62,7 +62,7 @@ __kernel void m11900_init (KERN_ATTR_TMPS_ESALT (pbkdf2_md5_tmp_t, pbkdf2_md5_t) md5_hmac_ctx_t md5_hmac_ctx; - md5_hmac_init_global (&md5_hmac_ctx, pws[gid].i, pws[gid].pw_len); + md5_hmac_init_global (&md5_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = md5_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = md5_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m12000-pure.cl b/OpenCL/m12000-pure.cl index e9ac20e59..a600c2b48 100644 --- a/OpenCL/m12000-pure.cl +++ b/OpenCL/m12000-pure.cl @@ -64,7 +64,7 @@ __kernel void m12000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, pbkdf2_sha1_ sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m12200-pure.cl b/OpenCL/m12200-pure.cl index 2f8ddc03f..2abde61b5 100644 --- a/OpenCL/m12200-pure.cl +++ b/OpenCL/m12200-pure.cl @@ -32,7 +32,7 @@ __kernel void m12200_init (KERN_ATTR_TMPS (ecryptfs_tmp_t)) sha512_update_global (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); - sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha512_final (&ctx); diff --git a/OpenCL/m12300-pure.cl b/OpenCL/m12300-pure.cl index 0dddcacf5..0a99d8fc9 100644 --- a/OpenCL/m12300-pure.cl +++ b/OpenCL/m12300-pure.cl @@ -86,7 +86,7 @@ __kernel void m12300_init (KERN_ATTR_TMPS (oraclet_tmp_t)) sha512_hmac_ctx_t sha512_hmac_ctx; - sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m12600_a1-pure.cl b/OpenCL/m12600_a1-pure.cl index ebde21a2e..c0940de6e 100644 --- a/OpenCL/m12600_a1-pure.cl +++ b/OpenCL/m12600_a1-pure.cl @@ -78,7 +78,7 @@ __kernel void m12600_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -224,7 +224,7 @@ __kernel void m12600_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m12700-pure.cl b/OpenCL/m12700-pure.cl index 07205c716..959004a1c 100644 --- a/OpenCL/m12700-pure.cl +++ b/OpenCL/m12700-pure.cl @@ -65,7 +65,7 @@ __kernel void m12700_init (KERN_ATTR_TMPS (mywallet_tmp_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m12800-pure.cl b/OpenCL/m12800-pure.cl index 059c3aba7..b14ddb1ac 100644 --- a/OpenCL/m12800-pure.cl +++ b/OpenCL/m12800-pure.cl @@ -96,7 +96,7 @@ __kernel void m12800_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha md4_init (&md4_ctx); - md4_update_global_utf16le (&md4_ctx, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&md4_ctx, pws[gid].i, pws[gid].pw_len & 255); md4_final (&md4_ctx); diff --git a/OpenCL/m12900-pure.cl b/OpenCL/m12900-pure.cl index a6999ceab..98f52a9d0 100644 --- a/OpenCL/m12900-pure.cl +++ b/OpenCL/m12900-pure.cl @@ -70,7 +70,7 @@ __kernel void m12900_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m13000-pure.cl b/OpenCL/m13000-pure.cl index b371def0c..b84f725aa 100644 --- a/OpenCL/m13000-pure.cl +++ b/OpenCL/m13000-pure.cl @@ -70,7 +70,7 @@ __kernel void m13000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m13100_a1-pure.cl b/OpenCL/m13100_a1-pure.cl index 5e0988168..26166dbb9 100644 --- a/OpenCL/m13100_a1-pure.cl +++ b/OpenCL/m13100_a1-pure.cl @@ -405,7 +405,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_A md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -465,7 +465,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_sxx (KERN_A md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m13200-pure.cl b/OpenCL/m13200-pure.cl index 567a823c6..6e88b22c8 100644 --- a/OpenCL/m13200-pure.cl +++ b/OpenCL/m13200-pure.cl @@ -29,7 +29,7 @@ __kernel void m13200_init (KERN_ATTR_TMPS (axcrypt_tmp_t)) sha1_init (&ctx); - sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha1_final (&ctx); diff --git a/OpenCL/m13300_a1-pure.cl b/OpenCL/m13300_a1-pure.cl index ede5b5ac9..6ce60d965 100644 --- a/OpenCL/m13300_a1-pure.cl +++ b/OpenCL/m13300_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m13300_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -88,7 +88,7 @@ __kernel void m13300_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m13400-pure.cl b/OpenCL/m13400-pure.cl index 5c207751a..25af48862 100644 --- a/OpenCL/m13400-pure.cl +++ b/OpenCL/m13400-pure.cl @@ -29,7 +29,7 @@ __kernel void m13400_init (KERN_ATTR_TMPS_ESALT (keepass_tmp_t, keepass_t)) sha256_init (&ctx); - sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha256_final (&ctx); diff --git a/OpenCL/m13500_a1-pure.cl b/OpenCL/m13500_a1-pure.cl index 526609e0e..37eb59adb 100644 --- a/OpenCL/m13500_a1-pure.cl +++ b/OpenCL/m13500_a1-pure.cl @@ -61,7 +61,7 @@ __kernel void m13500_mxx (KERN_ATTR_ESALT (pstoken_t)) * base */ - sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -144,7 +144,7 @@ __kernel void m13500_sxx (KERN_ATTR_ESALT (pstoken_t)) * base */ - sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m13600-pure.cl b/OpenCL/m13600-pure.cl index cb4f967d0..157ac14f9 100644 --- a/OpenCL/m13600-pure.cl +++ b/OpenCL/m13600-pure.cl @@ -64,7 +64,7 @@ __kernel void m13600_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, zip2_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m13800_a1-pure.cl b/OpenCL/m13800_a1-pure.cl index 9a3d73538..7db54e87e 100644 --- a/OpenCL/m13800_a1-pure.cl +++ b/OpenCL/m13800_a1-pure.cl @@ -32,7 +32,7 @@ __kernel void m13800_mxx (KERN_ATTR_ESALT (win8phone_t)) sha256_init (&ctx0); - sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -88,7 +88,7 @@ __kernel void m13800_sxx (KERN_ATTR_ESALT (win8phone_t)) sha256_init (&ctx0); - sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m13900_a1-pure.cl b/OpenCL/m13900_a1-pure.cl index ca0579742..791ec3d43 100644 --- a/OpenCL/m13900_a1-pure.cl +++ b/OpenCL/m13900_a1-pure.cl @@ -68,7 +68,7 @@ __kernel void m13900_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx1l); - sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -229,7 +229,7 @@ __kernel void m13900_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx1l); - sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len); + sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m14400_a1-pure.cl b/OpenCL/m14400_a1-pure.cl index ecce3323d..10ba89ce4 100644 --- a/OpenCL/m14400_a1-pure.cl +++ b/OpenCL/m14400_a1-pure.cl @@ -139,7 +139,7 @@ __kernel void m14400_mxx (KERN_ATTR_BASIC ()) sha1_update_64 (&ctx1, d20, d21, d22, d23, 2); - sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len); + 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); @@ -229,7 +229,7 @@ __kernel void m14400_mxx (KERN_ATTR_BASIC ()) sha1_update_64 (&ctx, d20, d21, d22, d23, 2); - sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + 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); @@ -396,7 +396,7 @@ __kernel void m14400_sxx (KERN_ATTR_BASIC ()) sha1_update_64 (&ctx1, d20, d21, d22, d23, 2); - sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len); + 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); @@ -486,7 +486,7 @@ __kernel void m14400_sxx (KERN_ATTR_BASIC ()) sha1_update_64 (&ctx, d20, d21, d22, d23, 2); - sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + 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); diff --git a/OpenCL/m14611-pure.cl b/OpenCL/m14611-pure.cl index 139dc29d2..689e7339e 100644 --- a/OpenCL/m14611-pure.cl +++ b/OpenCL/m14611-pure.cl @@ -73,7 +73,7 @@ __kernel void m14611_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14612-pure.cl b/OpenCL/m14612-pure.cl index d8db96cdb..a64a2a65d 100644 --- a/OpenCL/m14612-pure.cl +++ b/OpenCL/m14612-pure.cl @@ -73,7 +73,7 @@ __kernel void m14612_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14613-pure.cl b/OpenCL/m14613-pure.cl index c7315c6fe..824466943 100644 --- a/OpenCL/m14613-pure.cl +++ b/OpenCL/m14613-pure.cl @@ -73,7 +73,7 @@ __kernel void m14613_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14621-pure.cl b/OpenCL/m14621-pure.cl index ade2c2ff4..1ed24c4e9 100644 --- a/OpenCL/m14621-pure.cl +++ b/OpenCL/m14621-pure.cl @@ -79,7 +79,7 @@ __kernel void m14621_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14622-pure.cl b/OpenCL/m14622-pure.cl index 7594ca53c..9b64620a8 100644 --- a/OpenCL/m14622-pure.cl +++ b/OpenCL/m14622-pure.cl @@ -79,7 +79,7 @@ __kernel void m14622_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14623-pure.cl b/OpenCL/m14623-pure.cl index 1f70d96d3..fddf624ff 100644 --- a/OpenCL/m14623-pure.cl +++ b/OpenCL/m14623-pure.cl @@ -79,7 +79,7 @@ __kernel void m14623_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14631-pure.cl b/OpenCL/m14631-pure.cl index 94780f865..8b4b5f7c3 100644 --- a/OpenCL/m14631-pure.cl +++ b/OpenCL/m14631-pure.cl @@ -95,7 +95,7 @@ __kernel void m14631_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha512_hmac_ctx_t sha512_hmac_ctx; - sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad64[0] = sha512_hmac_ctx.ipad.h[0]; tmps[gid].ipad64[1] = sha512_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14632-pure.cl b/OpenCL/m14632-pure.cl index 24b647db0..1f49282c4 100644 --- a/OpenCL/m14632-pure.cl +++ b/OpenCL/m14632-pure.cl @@ -95,7 +95,7 @@ __kernel void m14632_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha512_hmac_ctx_t sha512_hmac_ctx; - sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad64[0] = sha512_hmac_ctx.ipad.h[0]; tmps[gid].ipad64[1] = sha512_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14633-pure.cl b/OpenCL/m14633-pure.cl index c12632dfc..07f033b4d 100644 --- a/OpenCL/m14633-pure.cl +++ b/OpenCL/m14633-pure.cl @@ -95,7 +95,7 @@ __kernel void m14633_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) sha512_hmac_ctx_t sha512_hmac_ctx; - sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad64[0] = sha512_hmac_ctx.ipad.h[0]; tmps[gid].ipad64[1] = sha512_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14641-pure.cl b/OpenCL/m14641-pure.cl index f1dd85508..a877f5b9f 100644 --- a/OpenCL/m14641-pure.cl +++ b/OpenCL/m14641-pure.cl @@ -73,7 +73,7 @@ __kernel void m14641_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) ripemd160_hmac_ctx_t ripemd160_hmac_ctx; - ripemd160_hmac_init_global (&ripemd160_hmac_ctx, pws[gid].i, pws[gid].pw_len); + ripemd160_hmac_init_global (&ripemd160_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = ripemd160_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = ripemd160_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14642-pure.cl b/OpenCL/m14642-pure.cl index 2fbbadb8b..e6c19251d 100644 --- a/OpenCL/m14642-pure.cl +++ b/OpenCL/m14642-pure.cl @@ -73,7 +73,7 @@ __kernel void m14642_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) ripemd160_hmac_ctx_t ripemd160_hmac_ctx; - ripemd160_hmac_init_global (&ripemd160_hmac_ctx, pws[gid].i, pws[gid].pw_len); + ripemd160_hmac_init_global (&ripemd160_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = ripemd160_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = ripemd160_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14643-pure.cl b/OpenCL/m14643-pure.cl index 23895de0e..cd0d22338 100644 --- a/OpenCL/m14643-pure.cl +++ b/OpenCL/m14643-pure.cl @@ -73,7 +73,7 @@ __kernel void m14643_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t)) ripemd160_hmac_ctx_t ripemd160_hmac_ctx; - ripemd160_hmac_init_global (&ripemd160_hmac_ctx, pws[gid].i, pws[gid].pw_len); + ripemd160_hmac_init_global (&ripemd160_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad32[0] = ripemd160_hmac_ctx.ipad.h[0]; tmps[gid].ipad32[1] = ripemd160_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14700-pure.cl b/OpenCL/m14700-pure.cl index 26d386245..bef7feeba 100644 --- a/OpenCL/m14700-pure.cl +++ b/OpenCL/m14700-pure.cl @@ -65,7 +65,7 @@ __kernel void m14700_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, itunes_backu sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m14800-pure.cl b/OpenCL/m14800-pure.cl index 09ed34dbd..7dce3fb3b 100644 --- a/OpenCL/m14800-pure.cl +++ b/OpenCL/m14800-pure.cl @@ -105,7 +105,7 @@ __kernel void m14800_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, itunes_bac sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m15000_a1-pure.cl b/OpenCL/m15000_a1-pure.cl index fd2ab69d4..88e6e3111 100644 --- a/OpenCL/m15000_a1-pure.cl +++ b/OpenCL/m15000_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m15000_mxx (KERN_ATTR_BASIC ()) sha512_init (&ctx0); - sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -106,7 +106,7 @@ __kernel void m15000_sxx (KERN_ATTR_BASIC ()) sha512_init (&ctx0); - sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m15100-pure.cl b/OpenCL/m15100-pure.cl index e79ffb5e2..b51c023f7 100644 --- a/OpenCL/m15100-pure.cl +++ b/OpenCL/m15100-pure.cl @@ -64,7 +64,7 @@ __kernel void m15100_init (KERN_ATTR_TMPS (pbkdf1_sha1_tmp_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m15300-pure.cl b/OpenCL/m15300-pure.cl index 4bc0ac397..79df1bc2b 100644 --- a/OpenCL/m15300-pure.cl +++ b/OpenCL/m15300-pure.cl @@ -602,7 +602,7 @@ __kernel void m15300_init (KERN_ATTR_TMPS_ESALT (dpapimk_tmp_v1_t, dpapimk_t)) sha1_init (&ctx); - sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha1_final (&ctx); @@ -620,7 +620,7 @@ __kernel void m15300_init (KERN_ATTR_TMPS_ESALT (dpapimk_tmp_v1_t, dpapimk_t)) md4_init (&ctx); - md4_update_global_utf16le (&ctx, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx, pws[gid].i, pws[gid].pw_len & 255); md4_final (&ctx); diff --git a/OpenCL/m15500_a1-pure.cl b/OpenCL/m15500_a1-pure.cl index 3883493fc..da1b33afe 100644 --- a/OpenCL/m15500_a1-pure.cl +++ b/OpenCL/m15500_a1-pure.cl @@ -41,7 +41,7 @@ __kernel void m15500_mxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_utf16be_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16be_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -112,7 +112,7 @@ __kernel void m15500_sxx (KERN_ATTR_BASIC ()) sha1_init (&ctx0); - sha1_update_global_utf16be_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16be_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m15600-pure.cl b/OpenCL/m15600-pure.cl index 711903a13..1cc52cb56 100644 --- a/OpenCL/m15600-pure.cl +++ b/OpenCL/m15600-pure.cl @@ -201,7 +201,7 @@ __kernel void m15600_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, ethereum_p sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m15700-pure.cl b/OpenCL/m15700-pure.cl index 8b2ff0b99..353ccd5e6 100644 --- a/OpenCL/m15700-pure.cl +++ b/OpenCL/m15700-pure.cl @@ -348,7 +348,7 @@ __kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_init (KERN_A sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len); @@ -456,7 +456,7 @@ __kernel void __attribute__((reqd_work_group_size(1, 1, 1))) m15700_comp (KERN_A sha256_hmac_ctx_t ctx; - sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); for (u32 l = 0; l < SCRYPT_CNT4; l += 4) { diff --git a/OpenCL/m15900-pure.cl b/OpenCL/m15900-pure.cl index b24ef959e..4a9d8d3a1 100644 --- a/OpenCL/m15900-pure.cl +++ b/OpenCL/m15900-pure.cl @@ -101,7 +101,7 @@ __kernel void m15900_init (KERN_ATTR_TMPS_ESALT (dpapimk_tmp_v2_t, dpapimk_t)) sha1_init (&ctx); - sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len); + sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255); sha1_final (&ctx); @@ -119,7 +119,7 @@ __kernel void m15900_init (KERN_ATTR_TMPS_ESALT (dpapimk_tmp_v2_t, dpapimk_t)) md4_init (&ctx); - md4_update_global_utf16le (&ctx, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx, pws[gid].i, pws[gid].pw_len & 255); md4_final (&ctx); diff --git a/OpenCL/m16100_a1-pure.cl b/OpenCL/m16100_a1-pure.cl index 749db59aa..68f7dadbc 100644 --- a/OpenCL/m16100_a1-pure.cl +++ b/OpenCL/m16100_a1-pure.cl @@ -55,7 +55,7 @@ __kernel void m16100_mxx (KERN_ATTR_ESALT (tacacs_plus_t)) md5_update_64 (&ctx0, session0, session1, session2, session3, 4); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); u32 ct_buf[2]; @@ -213,7 +213,7 @@ __kernel void m16100_sxx (KERN_ATTR_ESALT (tacacs_plus_t)) md5_update_64 (&ctx0, session0, session1, session2, session3, 4); - md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); u32 ct_buf[2]; diff --git a/OpenCL/m16200-pure.cl b/OpenCL/m16200-pure.cl index 51e944509..5dc1f72f0 100644 --- a/OpenCL/m16200-pure.cl +++ b/OpenCL/m16200-pure.cl @@ -71,7 +71,7 @@ __kernel void m16200_init (KERN_ATTR_TMPS_ESALT (apple_secure_notes_tmp_t, apple sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m16300-pure.cl b/OpenCL/m16300-pure.cl index 8ffb4129e..dac913354 100644 --- a/OpenCL/m16300-pure.cl +++ b/OpenCL/m16300-pure.cl @@ -202,7 +202,7 @@ __kernel void m16300_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, ethereum_p sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; @@ -222,7 +222,7 @@ __kernel void m16300_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, ethereum_p tmps[gid].opad[6] = sha256_hmac_ctx.opad.h[6]; tmps[gid].opad[7] = sha256_hmac_ctx.opad.h[7]; - sha256_hmac_update_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_update_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); for (u32 i = 0, j = 1; i < 8; i += 8, j += 1) { diff --git a/OpenCL/m16400_a1-pure.cl b/OpenCL/m16400_a1-pure.cl index 6e929dfc3..366d77075 100644 --- a/OpenCL/m16400_a1-pure.cl +++ b/OpenCL/m16400_a1-pure.cl @@ -194,7 +194,7 @@ __kernel void m16400_mxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - cram_md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + cram_md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -248,7 +248,7 @@ __kernel void m16400_sxx (KERN_ATTR_BASIC ()) md5_init (&ctx0); - cram_md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len); + cram_md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m16600_a1-pure.cl b/OpenCL/m16600_a1-pure.cl index c8d056b4c..69bd4e2d5 100644 --- a/OpenCL/m16600_a1-pure.cl +++ b/OpenCL/m16600_a1-pure.cl @@ -85,7 +85,7 @@ __kernel void m16600_mxx (KERN_ATTR_ESALT (electrum_wallet_t)) sha256_init (&ctx0); - sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -258,7 +258,7 @@ __kernel void m16600_sxx (KERN_ATTR_ESALT (electrum_wallet_t)) sha256_init (&ctx0); - sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len); + sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m16800-pure.cl b/OpenCL/m16800-pure.cl index 5526d4250..a44e72940 100644 --- a/OpenCL/m16800-pure.cl +++ b/OpenCL/m16800-pure.cl @@ -64,7 +64,7 @@ __kernel void m16800_init (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_pmkid_t)) sha1_hmac_ctx_t sha1_hmac_ctx; - sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m16900-pure.cl b/OpenCL/m16900-pure.cl index 89322a47a..f879a8d4a 100644 --- a/OpenCL/m16900-pure.cl +++ b/OpenCL/m16900-pure.cl @@ -70,7 +70,7 @@ __kernel void m16900_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, ansible_va sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1]; diff --git a/OpenCL/m18200_a1-pure.cl b/OpenCL/m18200_a1-pure.cl index dc6d6d9f1..a1535dfe4 100644 --- a/OpenCL/m18200_a1-pure.cl +++ b/OpenCL/m18200_a1-pure.cl @@ -403,7 +403,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_A md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop @@ -463,7 +463,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_sxx (KERN_A md4_init (&ctx0); - md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len); + md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255); /** * loop diff --git a/OpenCL/m18300-pure.cl b/OpenCL/m18300-pure.cl index b8c5ed288..86a117d44 100644 --- a/OpenCL/m18300-pure.cl +++ b/OpenCL/m18300-pure.cl @@ -71,7 +71,7 @@ __kernel void m18300_init (KERN_ATTR_TMPS_ESALT (apple_secure_notes_tmp_t, apple sha256_hmac_ctx_t sha256_hmac_ctx; - sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len); + sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255); tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0]; tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];