arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-29 11:28:15 +00:00
Code Issues Releases Wiki Activity

Give the compiler a hint for automatic optimizations based on password length

Browse Source
This commit is contained in:
Jens Steube 2018-11-20 15:26:46 +01:00
parent 490050ecb3
commit 2a6444c05a
363 changed files with 1124 additions and 1124 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
OpenCL/m00000_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00000_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -177,7 +177,7 @@ __kernel void m00000_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m00000_a3-optimized.cl
View File

@ -457,7 +457,7 @@ __kernel void m00000_m04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -495,7 +495,7 @@ __kernel void m00000_m08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -533,7 +533,7 @@ __kernel void m00000_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -571,7 +571,7 @@ __kernel void m00000_s04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -609,7 +609,7 @@ __kernel void m00000_s08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -647,7 +647,7 @@ __kernel void m00000_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00000_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00000_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -95,7 +95,7 @@ __kernel void m00000_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00010_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00010_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -251,7 +251,7 @@ __kernel void m00010_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m00010_a3-optimized.cl
View File

@ -505,7 +505,7 @@ __kernel void m00010_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -543,7 +543,7 @@ __kernel void m00010_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -581,7 +581,7 @@ __kernel void m00010_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -619,7 +619,7 @@ __kernel void m00010_s04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -657,7 +657,7 @@ __kernel void m00010_s08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -695,7 +695,7 @@ __kernel void m00010_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00010_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00010_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m00010_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00020_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00020_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -231,7 +231,7 @@ __kernel void m00020_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m00020_a3-optimized.cl
View File

@ -415,7 +415,7 @@ __kernel void m00020_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -462,7 +462,7 @@ __kernel void m00020_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -509,7 +509,7 @@ __kernel void m00020_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -556,7 +556,7 @@ __kernel void m00020_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -603,7 +603,7 @@ __kernel void m00020_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -650,7 +650,7 @@ __kernel void m00020_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00020_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00020_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -101,7 +101,7 @@ __kernel void m00020_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00030_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00030_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -256,7 +256,7 @@ __kernel void m00030_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m00030_a3-optimized.cl
View File

@ -505,7 +505,7 @@ __kernel void m00030_m04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -543,7 +543,7 @@ __kernel void m00030_m08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -581,7 +581,7 @@ __kernel void m00030_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -619,7 +619,7 @@ __kernel void m00030_s04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -657,7 +657,7 @@ __kernel void m00030_s08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -695,7 +695,7 @@ __kernel void m00030_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00030_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00030_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m00030_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00040_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00040_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -236,7 +236,7 @@ __kernel void m00040_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m00040_a3-optimized.cl
View File

@ -415,7 +415,7 @@ __kernel void m00040_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -462,7 +462,7 @@ __kernel void m00040_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -509,7 +509,7 @@ __kernel void m00040_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -556,7 +556,7 @@ __kernel void m00040_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -603,7 +603,7 @@ __kernel void m00040_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -650,7 +650,7 @@ __kernel void m00040_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00040_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00040_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -101,7 +101,7 @@ __kernel void m00040_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00050_a0-optimized.cl
View File

@ -128,7 +128,7 @@ __kernel void m00050_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -241,7 +241,7 @@ __kernel void m00050_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

4
OpenCL/m00050_a1-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00050_mxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };
@ -118,7 +118,7 @@ __kernel void m00050_sxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

12
OpenCL/m00050_a3-optimized.cl
View File

@ -358,7 +358,7 @@ __kernel void m00050_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -405,7 +405,7 @@ __kernel void m00050_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -452,7 +452,7 @@ __kernel void m00050_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -499,7 +499,7 @@ __kernel void m00050_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -546,7 +546,7 @@ __kernel void m00050_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -593,7 +593,7 @@ __kernel void m00050_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00050_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00050_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -104,7 +104,7 @@ __kernel void m00050_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00060_a0-optimized.cl
View File

@ -128,7 +128,7 @@ __kernel void m00060_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -265,7 +265,7 @@ __kernel void m00060_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

4
OpenCL/m00060_a1-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00060_mxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };
@ -120,7 +120,7 @@ __kernel void m00060_sxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

12
OpenCL/m00060_a3-optimized.cl
View File

@ -358,7 +358,7 @@ __kernel void m00060_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -405,7 +405,7 @@ __kernel void m00060_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -452,7 +452,7 @@ __kernel void m00060_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -499,7 +499,7 @@ __kernel void m00060_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -546,7 +546,7 @@ __kernel void m00060_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -593,7 +593,7 @@ __kernel void m00060_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00060_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00060_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m00060_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00100_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00100_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -222,7 +222,7 @@ __kernel void m00100_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m00100_a3-optimized.cl
View File

@ -549,7 +549,7 @@ __kernel void m00100_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -587,7 +587,7 @@ __kernel void m00100_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -625,7 +625,7 @@ __kernel void m00100_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -663,7 +663,7 @@ __kernel void m00100_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -701,7 +701,7 @@ __kernel void m00100_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -739,7 +739,7 @@ __kernel void m00100_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00100_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00100_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -95,7 +95,7 @@ __kernel void m00100_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00110_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00110_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -295,7 +295,7 @@ __kernel void m00110_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m00110_a3-optimized.cl
View File

@ -597,7 +597,7 @@ __kernel void m00110_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -635,7 +635,7 @@ __kernel void m00110_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -673,7 +673,7 @@ __kernel void m00110_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -711,7 +711,7 @@ __kernel void m00110_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -749,7 +749,7 @@ __kernel void m00110_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -787,7 +787,7 @@ __kernel void m00110_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00110_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00110_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m00110_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00120_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00120_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -287,7 +287,7 @@ __kernel void m00120_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m00120_a3-optimized.cl
View File

@ -509,7 +509,7 @@ __kernel void m00120_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -556,7 +556,7 @@ __kernel void m00120_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -603,7 +603,7 @@ __kernel void m00120_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -650,7 +650,7 @@ __kernel void m00120_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -697,7 +697,7 @@ __kernel void m00120_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -744,7 +744,7 @@ __kernel void m00120_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00120_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00120_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -101,7 +101,7 @@ __kernel void m00120_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00130_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00130_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -300,7 +300,7 @@ __kernel void m00130_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m00130_a3-optimized.cl
View File

@ -597,7 +597,7 @@ __kernel void m00130_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -635,7 +635,7 @@ __kernel void m00130_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -673,7 +673,7 @@ __kernel void m00130_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -711,7 +711,7 @@ __kernel void m00130_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -749,7 +749,7 @@ __kernel void m00130_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -787,7 +787,7 @@ __kernel void m00130_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00130_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00130_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m00130_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00140_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00140_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -280,7 +280,7 @@ __kernel void m00140_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m00140_a3-optimized.cl
View File

@ -509,7 +509,7 @@ __kernel void m00140_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -556,7 +556,7 @@ __kernel void m00140_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -603,7 +603,7 @@ __kernel void m00140_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -650,7 +650,7 @@ __kernel void m00140_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -697,7 +697,7 @@ __kernel void m00140_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -744,7 +744,7 @@ __kernel void m00140_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00140_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00140_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -101,7 +101,7 @@ __kernel void m00140_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00150_a0-optimized.cl
View File

@ -132,7 +132,7 @@ __kernel void m00150_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -254,7 +254,7 @@ __kernel void m00150_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

4
OpenCL/m00150_a1-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00150_mxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };
@ -118,7 +118,7 @@ __kernel void m00150_sxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

12
OpenCL/m00150_a3-optimized.cl
View File

@ -362,7 +362,7 @@ __kernel void m00150_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -409,7 +409,7 @@ __kernel void m00150_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -456,7 +456,7 @@ __kernel void m00150_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -503,7 +503,7 @@ __kernel void m00150_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -550,7 +550,7 @@ __kernel void m00150_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -597,7 +597,7 @@ __kernel void m00150_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00150_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00150_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -104,7 +104,7 @@ __kernel void m00150_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00160_a0-optimized.cl
View File

@ -132,7 +132,7 @@ __kernel void m00160_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -269,7 +269,7 @@ __kernel void m00160_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

4
OpenCL/m00160_a1-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00160_mxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };
@ -120,7 +120,7 @@ __kernel void m00160_sxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

12
OpenCL/m00160_a3-optimized.cl
View File

@ -362,7 +362,7 @@ __kernel void m00160_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -409,7 +409,7 @@ __kernel void m00160_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -456,7 +456,7 @@ __kernel void m00160_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -503,7 +503,7 @@ __kernel void m00160_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -550,7 +550,7 @@ __kernel void m00160_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -597,7 +597,7 @@ __kernel void m00160_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00160_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00160_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m00160_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m00200_a0-optimized.cl
View File

@ -43,7 +43,7 @@ __kernel void m00200_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -169,7 +169,7 @@ __kernel void m00200_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m00200_a3-optimized.cl
View File

@ -422,7 +422,7 @@ __kernel void m00200_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -460,7 +460,7 @@ __kernel void m00200_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -498,7 +498,7 @@ __kernel void m00200_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -536,7 +536,7 @@ __kernel void m00200_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -574,7 +574,7 @@ __kernel void m00200_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -612,7 +612,7 @@ __kernel void m00200_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00300_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00300_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -347,7 +347,7 @@ __kernel void m00300_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m00300_a3-optimized.cl
View File

@ -800,7 +800,7 @@ __kernel void m00300_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -838,7 +838,7 @@ __kernel void m00300_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -876,7 +876,7 @@ __kernel void m00300_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -914,7 +914,7 @@ __kernel void m00300_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -952,7 +952,7 @@ __kernel void m00300_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -990,7 +990,7 @@ __kernel void m00300_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00300_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00300_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -118,7 +118,7 @@ __kernel void m00300_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

2
OpenCL/m00400-optimized.cl
View File

@ -47,7 +47,7 @@ __kernel void m00400_init (KERN_ATTR_TMPS (phpass_tmp_t))
w2[2] = 0;
w2[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

2
OpenCL/m00400-pure.cl
View File

@ -67,7 +67,7 @@ __kernel void m00400_loop (KERN_ATTR_TMPS (phpass_tmp_t))
* init
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

4
OpenCL/m00500-optimized.cl
View File

@ -667,7 +667,7 @@ __kernel void m00500_init (KERN_ATTR_TMPS (md5crypt_tmp_t))
w0[2] = pws[gid].i[2];
w0[3] = pws[gid].i[3];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -830,7 +830,7 @@ __kernel void m00500_loop (KERN_ATTR_TMPS (md5crypt_tmp_t))
w0[2] = pws[gid].i[2];
w0[3] = pws[gid].i[3];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
u32 w0_x80[4];

4
OpenCL/m00500-pure.cl
View File

@ -31,7 +31,7 @@ __kernel void m00500_init (KERN_ATTR_TMPS (md5crypt_tmp_t))
* init
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };
@ -137,7 +137,7 @@ __kernel void m00500_loop (KERN_ATTR_TMPS (md5crypt_tmp_t))
* init
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

4
OpenCL/m00600_a0-optimized.cl
View File

@ -142,7 +142,7 @@ __kernel void m00600_m04 (KERN_ATTR_RULES_ESALT (blake2_t))
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
u64 tmp_h[8];
u64 tmp_t[2];
@ -250,7 +250,7 @@ __kernel void m00600_s04 (KERN_ATTR_RULES_ESALT (blake2_t))
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
u64 tmp_h[8];
u64 tmp_t[2];

4
OpenCL/m00900_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m00900_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -158,7 +158,7 @@ __kernel void m00900_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m00900_a3-optimized.cl
View File

@ -380,7 +380,7 @@ __kernel void m00900_m04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -418,7 +418,7 @@ __kernel void m00900_m08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -456,7 +456,7 @@ __kernel void m00900_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -494,7 +494,7 @@ __kernel void m00900_s04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -532,7 +532,7 @@ __kernel void m00900_s08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -570,7 +570,7 @@ __kernel void m00900_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m00900_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m00900_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -95,7 +95,7 @@ __kernel void m00900_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01000_a0-optimized.cl
View File

@ -42,7 +42,7 @@ __kernel void m01000_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -161,7 +161,7 @@ __kernel void m01000_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m01000_a3-optimized.cl
View File

@ -380,7 +380,7 @@ __kernel void m01000_m04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -418,7 +418,7 @@ __kernel void m01000_m08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -456,7 +456,7 @@ __kernel void m01000_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -494,7 +494,7 @@ __kernel void m01000_s04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -532,7 +532,7 @@ __kernel void m01000_s08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -570,7 +570,7 @@ __kernel void m01000_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01000_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01000_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -95,7 +95,7 @@ __kernel void m01000_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01100_a0-optimized.cl
View File

@ -35,7 +35,7 @@ __kernel void m01100_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -254,7 +254,7 @@ __kernel void m01100_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

8
OpenCL/m01100_a3-optimized.cl
View File

@ -525,7 +525,7 @@ __kernel void m01100_m04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -579,7 +579,7 @@ __kernel void m01100_m08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -637,7 +637,7 @@ __kernel void m01100_s04 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -691,7 +691,7 @@ __kernel void m01100_s08 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

4
OpenCL/m01100_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01100_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -119,7 +119,7 @@ __kernel void m01100_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01300_a0-optimized.cl
View File

@ -55,7 +55,7 @@ __kernel void m01300_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -210,7 +210,7 @@ __kernel void m01300_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m01300_a3-optimized.cl
View File

@ -323,7 +323,7 @@ __kernel void m01300_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -361,7 +361,7 @@ __kernel void m01300_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -399,7 +399,7 @@ __kernel void m01300_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -437,7 +437,7 @@ __kernel void m01300_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -475,7 +475,7 @@ __kernel void m01300_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -513,7 +513,7 @@ __kernel void m01300_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01300_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01300_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -95,7 +95,7 @@ __kernel void m01300_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01400_a0-optimized.cl
View File

@ -56,7 +56,7 @@ __kernel void m01400_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -211,7 +211,7 @@ __kernel void m01400_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m01400_a3-optimized.cl
View File

@ -326,7 +326,7 @@ __kernel void m01400_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -364,7 +364,7 @@ __kernel void m01400_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -402,7 +402,7 @@ __kernel void m01400_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -440,7 +440,7 @@ __kernel void m01400_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -478,7 +478,7 @@ __kernel void m01400_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -516,7 +516,7 @@ __kernel void m01400_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01400_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01400_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -95,7 +95,7 @@ __kernel void m01400_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01410_a0-optimized.cl
View File

@ -56,7 +56,7 @@ __kernel void m01410_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -284,7 +284,7 @@ __kernel void m01410_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m01410_a3-optimized.cl
View File

@ -377,7 +377,7 @@ __kernel void m01410_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -415,7 +415,7 @@ __kernel void m01410_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -453,7 +453,7 @@ __kernel void m01410_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -491,7 +491,7 @@ __kernel void m01410_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -529,7 +529,7 @@ __kernel void m01410_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -567,7 +567,7 @@ __kernel void m01410_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01410_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01410_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m01410_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01420_a0-optimized.cl
View File

@ -56,7 +56,7 @@ __kernel void m01420_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -264,7 +264,7 @@ __kernel void m01420_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m01420_a3-optimized.cl
View File

@ -485,7 +485,7 @@ __kernel void m01420_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -532,7 +532,7 @@ __kernel void m01420_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -579,7 +579,7 @@ __kernel void m01420_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -626,7 +626,7 @@ __kernel void m01420_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -673,7 +673,7 @@ __kernel void m01420_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -720,7 +720,7 @@ __kernel void m01420_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01420_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01420_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -101,7 +101,7 @@ __kernel void m01420_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01430_a0-optimized.cl
View File

@ -56,7 +56,7 @@ __kernel void m01430_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -289,7 +289,7 @@ __kernel void m01430_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m01430_a3-optimized.cl
View File

@ -377,7 +377,7 @@ __kernel void m01430_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -415,7 +415,7 @@ __kernel void m01430_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -453,7 +453,7 @@ __kernel void m01430_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -491,7 +491,7 @@ __kernel void m01430_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -529,7 +529,7 @@ __kernel void m01430_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -567,7 +567,7 @@ __kernel void m01430_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01430_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01430_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m01430_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01440_a0-optimized.cl
View File

@ -56,7 +56,7 @@ __kernel void m01440_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -269,7 +269,7 @@ __kernel void m01440_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m01440_a3-optimized.cl
View File

@ -485,7 +485,7 @@ __kernel void m01440_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -532,7 +532,7 @@ __kernel void m01440_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -579,7 +579,7 @@ __kernel void m01440_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -626,7 +626,7 @@ __kernel void m01440_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -673,7 +673,7 @@ __kernel void m01440_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -720,7 +720,7 @@ __kernel void m01440_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01440_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01440_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -101,7 +101,7 @@ __kernel void m01440_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01450_a0-optimized.cl
View File

@ -144,7 +144,7 @@ __kernel void m01450_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -266,7 +266,7 @@ __kernel void m01450_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

4
OpenCL/m01450_a1-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01450_mxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };
@ -118,7 +118,7 @@ __kernel void m01450_sxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

12
OpenCL/m01450_a3-optimized.cl
View File

@ -374,7 +374,7 @@ __kernel void m01450_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -421,7 +421,7 @@ __kernel void m01450_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -468,7 +468,7 @@ __kernel void m01450_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -515,7 +515,7 @@ __kernel void m01450_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -562,7 +562,7 @@ __kernel void m01450_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -609,7 +609,7 @@ __kernel void m01450_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01450_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01450_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -104,7 +104,7 @@ __kernel void m01450_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01460_a0-optimized.cl
View File

@ -144,7 +144,7 @@ __kernel void m01460_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -281,7 +281,7 @@ __kernel void m01460_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

4
OpenCL/m01460_a1-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01460_mxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };
@ -120,7 +120,7 @@ __kernel void m01460_sxx (KERN_ATTR_BASIC ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

12
OpenCL/m01460_a3-optimized.cl
View File

@ -374,7 +374,7 @@ __kernel void m01460_m04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -421,7 +421,7 @@ __kernel void m01460_m08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -468,7 +468,7 @@ __kernel void m01460_m16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -515,7 +515,7 @@ __kernel void m01460_s04 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -562,7 +562,7 @@ __kernel void m01460_s08 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -609,7 +609,7 @@ __kernel void m01460_s16 (KERN_ATTR_BASIC ())
w3[2] = 0;
w3[3] = 0;
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01460_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01460_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m01460_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01600-optimized.cl
View File

@ -682,7 +682,7 @@ __kernel void m01600_init (KERN_ATTR_TMPS (md5crypt_tmp_t))
w0[2] = pws[gid].i[2];
w0[3] = pws[gid].i[3];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -845,7 +845,7 @@ __kernel void m01600_loop (KERN_ATTR_TMPS (md5crypt_tmp_t))
w0[2] = pws[gid].i[2];
w0[3] = pws[gid].i[3];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
u32 w0_x80[4];

4
OpenCL/m01600-pure.cl
View File

@ -35,7 +35,7 @@ __kernel void m01600_init (KERN_ATTR_TMPS (md5crypt_tmp_t))
* init
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };
@ -142,7 +142,7 @@ __kernel void m01600_loop (KERN_ATTR_TMPS (md5crypt_tmp_t))
* init
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32 w[64] = { 0 };

4
OpenCL/m01700_a0-optimized.cl
View File

@ -165,7 +165,7 @@ __kernel void m01700_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* loop
@ -266,7 +266,7 @@ __kernel void m01700_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* digest

12
OpenCL/m01700_a3-optimized.cl
View File

@ -306,7 +306,7 @@ __kernel void m01700_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -344,7 +344,7 @@ __kernel void m01700_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -382,7 +382,7 @@ __kernel void m01700_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -420,7 +420,7 @@ __kernel void m01700_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -458,7 +458,7 @@ __kernel void m01700_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -496,7 +496,7 @@ __kernel void m01700_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01700_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01700_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -95,7 +95,7 @@ __kernel void m01700_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01710_a0-optimized.cl
View File

@ -165,7 +165,7 @@ __kernel void m01710_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -339,7 +339,7 @@ __kernel void m01710_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

12
OpenCL/m01710_a3-optimized.cl
View File

@ -357,7 +357,7 @@ __kernel void m01710_m04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -395,7 +395,7 @@ __kernel void m01710_m08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -433,7 +433,7 @@ __kernel void m01710_m16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -471,7 +471,7 @@ __kernel void m01710_s04 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -509,7 +509,7 @@ __kernel void m01710_s08 (KERN_ATTR_VECTOR ())
w[14] = 0;
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main
@ -547,7 +547,7 @@ __kernel void m01710_s16 (KERN_ATTR_VECTOR ())
w[14] = pws[gid].i[14];
w[15] = pws[gid].i[15];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* main

4
OpenCL/m01710_a3-pure.cl
View File

@ -28,7 +28,7 @@ __kernel void m01710_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };
@ -106,7 +106,7 @@ __kernel void m01710_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 255;
u32x w[64] = { 0 };

4
OpenCL/m01720_a0-optimized.cl
View File

@ -165,7 +165,7 @@ __kernel void m01720_m04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt
@ -319,7 +319,7 @@ __kernel void m01720_s04 (KERN_ATTR_RULES ())
pw_buf1[2] = pws[gid].i[6];
pw_buf1[3] = pws[gid].i[7];
const u32 pw_len = pws[gid].pw_len;
const u32 pw_len = pws[gid].pw_len & 63;
/**
* salt

Some files were not shown because too many files have changed in this diff Show More