From 64dfd40113452e0601cfe5805aa7241932282767 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Tue, 20 Nov 2018 15:44:24 +0100 Subject: [PATCH] Give the compiler a hint for automatic optimizations based on password length --- OpenCL/m00000_a1-optimized.cl | 12 ++++++------ OpenCL/m00010_a1-optimized.cl | 12 ++++++------ OpenCL/m00020_a1-optimized.cl | 12 ++++++------ OpenCL/m00030_a1-optimized.cl | 12 ++++++------ OpenCL/m00040_a1-optimized.cl | 12 ++++++------ OpenCL/m00050_a1-optimized.cl | 12 ++++++------ OpenCL/m00060_a1-optimized.cl | 12 ++++++------ OpenCL/m00100_a1-optimized.cl | 12 ++++++------ OpenCL/m00110_a1-optimized.cl | 12 ++++++------ OpenCL/m00120_a1-optimized.cl | 12 ++++++------ OpenCL/m00130_a1-optimized.cl | 12 ++++++------ OpenCL/m00140_a1-optimized.cl | 12 ++++++------ OpenCL/m00150_a1-optimized.cl | 12 ++++++------ OpenCL/m00160_a1-optimized.cl | 12 ++++++------ OpenCL/m00200_a1-optimized.cl | 12 ++++++------ OpenCL/m00300_a1-optimized.cl | 12 ++++++------ OpenCL/m00600_a1-optimized.cl | 8 ++++---- OpenCL/m00900_a1-optimized.cl | 12 ++++++------ OpenCL/m01000_a1-optimized.cl | 12 ++++++------ OpenCL/m01100_a1-optimized.cl | 12 ++++++------ OpenCL/m01300_a1-optimized.cl | 12 ++++++------ OpenCL/m01400_a1-optimized.cl | 12 ++++++------ OpenCL/m01410_a1-optimized.cl | 12 ++++++------ OpenCL/m01420_a1-optimized.cl | 12 ++++++------ OpenCL/m01430_a1-optimized.cl | 12 ++++++------ OpenCL/m01440_a1-optimized.cl | 12 ++++++------ OpenCL/m01450_a1-optimized.cl | 12 ++++++------ OpenCL/m01460_a1-optimized.cl | 12 ++++++------ OpenCL/m01500_a1-pure.cl | 12 ++++++------ OpenCL/m01700_a1-optimized.cl | 12 ++++++------ OpenCL/m01710_a1-optimized.cl | 12 ++++++------ OpenCL/m01720_a1-optimized.cl | 12 ++++++------ OpenCL/m01730_a1-optimized.cl | 12 ++++++------ OpenCL/m01740_a1-optimized.cl | 12 ++++++------ OpenCL/m01750_a1-optimized.cl | 12 ++++++------ OpenCL/m01760_a1-optimized.cl | 12 ++++++------ OpenCL/m02400_a1-optimized.cl | 12 ++++++------ OpenCL/m02410_a1-optimized.cl | 12 ++++++------ OpenCL/m02610_a1-optimized.cl | 12 ++++++------ OpenCL/m02710_a1-optimized.cl | 12 ++++++------ OpenCL/m02810_a1-optimized.cl | 12 ++++++------ OpenCL/m03000_a1-pure.cl | 12 ++++++------ OpenCL/m03100_a1-optimized.cl | 12 ++++++------ OpenCL/m03710_a1-optimized.cl | 12 ++++++------ OpenCL/m03800_a1-optimized.cl | 12 ++++++------ OpenCL/m03910_a1-optimized.cl | 12 ++++++------ OpenCL/m04010_a1-optimized.cl | 12 ++++++------ OpenCL/m04110_a1-optimized.cl | 12 ++++++------ OpenCL/m04310_a1-optimized.cl | 12 ++++++------ OpenCL/m04400_a1-optimized.cl | 12 ++++++------ OpenCL/m04500_a1-optimized.cl | 12 ++++++------ OpenCL/m04520_a1-optimized.cl | 12 ++++++------ OpenCL/m04700_a1-optimized.cl | 12 ++++++------ OpenCL/m04800_a1-optimized.cl | 12 ++++++------ OpenCL/m04900_a1-optimized.cl | 12 ++++++------ OpenCL/m05100_a1-optimized.cl | 12 ++++++------ OpenCL/m05300_a1-optimized.cl | 12 ++++++------ OpenCL/m05400_a1-optimized.cl | 12 ++++++------ OpenCL/m05500_a1-optimized.cl | 12 ++++++------ OpenCL/m05600_a1-optimized.cl | 12 ++++++------ OpenCL/m06000_a1-optimized.cl | 12 ++++++------ OpenCL/m06100_a1-optimized.cl | 12 ++++++------ OpenCL/m06900_a1-optimized.cl | 12 ++++++------ OpenCL/m07000_a1-optimized.cl | 12 ++++++------ OpenCL/m07300_a1-optimized.cl | 12 ++++++------ OpenCL/m07500_a1-optimized.cl | 12 ++++++------ OpenCL/m07700_a1-optimized.cl | 12 ++++++------ OpenCL/m07701_a1-optimized.cl | 12 ++++++------ OpenCL/m07800_a1-optimized.cl | 12 ++++++------ OpenCL/m07801_a1-optimized.cl | 12 ++++++------ OpenCL/m08000_a1-optimized.cl | 12 ++++++------ OpenCL/m08100_a1-optimized.cl | 12 ++++++------ OpenCL/m08300_a1-optimized.cl | 12 ++++++------ OpenCL/m08400_a1-optimized.cl | 12 ++++++------ OpenCL/m08500_a1-pure.cl | 12 ++++++------ OpenCL/m08600_a1-pure.cl | 12 ++++++------ OpenCL/m08700_a1-optimized.cl | 12 ++++++------ OpenCL/m09700_a1-optimized.cl | 12 ++++++------ OpenCL/m09710_a1-optimized.cl | 12 ++++++------ OpenCL/m09720_a1-optimized.cl | 12 ++++++------ OpenCL/m09800_a1-optimized.cl | 12 ++++++------ OpenCL/m09810_a1-optimized.cl | 12 ++++++------ OpenCL/m09820_a1-optimized.cl | 12 ++++++------ OpenCL/m09900_a1-optimized.cl | 12 ++++++------ OpenCL/m10100_a1-optimized.cl | 12 ++++++------ OpenCL/m10400_a1-optimized.cl | 12 ++++++------ OpenCL/m10410_a1-optimized.cl | 12 ++++++------ OpenCL/m10420_a1-optimized.cl | 12 ++++++------ OpenCL/m10800_a1-optimized.cl | 12 ++++++------ OpenCL/m11000_a1-optimized.cl | 12 ++++++------ OpenCL/m11100_a1-optimized.cl | 12 ++++++------ OpenCL/m11200_a1-optimized.cl | 12 ++++++------ OpenCL/m11500_a1-optimized.cl | 12 ++++++------ OpenCL/m11700_a1-optimized.cl | 12 ++++++------ OpenCL/m11800_a1-optimized.cl | 12 ++++++------ OpenCL/m12600_a1-optimized.cl | 12 ++++++------ OpenCL/m13100_a1-optimized.cl | 12 ++++++------ OpenCL/m13300_a1-optimized.cl | 12 ++++++------ OpenCL/m13500_a1-optimized.cl | 12 ++++++------ OpenCL/m13800_a1-optimized.cl | 12 ++++++------ OpenCL/m13900_a1-optimized.cl | 12 ++++++------ OpenCL/m14000_a1-pure.cl | 12 ++++++------ OpenCL/m14100_a1-pure.cl | 12 ++++++------ OpenCL/m14400_a1-optimized.cl | 12 ++++++------ OpenCL/m14900_a1-optimized.cl | 12 ++++++------ OpenCL/m15000_a1-optimized.cl | 12 ++++++------ OpenCL/m15400_a1-optimized.cl | 8 ++++---- OpenCL/m15500_a1-optimized.cl | 12 ++++++------ OpenCL/m16000_a1-pure.cl | 12 ++++++------ OpenCL/m16100_a1-optimized.cl | 12 ++++++------ OpenCL/m16400_a1-optimized.cl | 12 ++++++------ OpenCL/m16600_a1-optimized.cl | 12 ++++++------ OpenCL/m17300_a1-optimized.cl | 12 ++++++------ OpenCL/m17400_a1-optimized.cl | 12 ++++++------ OpenCL/m17500_a1-optimized.cl | 12 ++++++------ OpenCL/m17600_a1-optimized.cl | 12 ++++++------ OpenCL/m17700_a1-optimized.cl | 12 ++++++------ OpenCL/m17800_a1-optimized.cl | 12 ++++++------ OpenCL/m17900_a1-optimized.cl | 12 ++++++------ OpenCL/m18000_a1-optimized.cl | 12 ++++++------ OpenCL/m18200_a1-optimized.cl | 12 ++++++------ 121 files changed, 722 insertions(+), 722 deletions(-) diff --git a/OpenCL/m00000_a1-optimized.cl b/OpenCL/m00000_a1-optimized.cl index a72976897..2bbd9324b 100644 --- a/OpenCL/m00000_a1-optimized.cl +++ b/OpenCL/m00000_a1-optimized.cl @@ -42,7 +42,7 @@ __kernel void m00000_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -50,9 +50,9 @@ __kernel void m00000_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -236,7 +236,7 @@ __kernel void m00000_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -256,9 +256,9 @@ __kernel void m00000_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00010_a1-optimized.cl b/OpenCL/m00010_a1-optimized.cl index f866b170f..1f4003f21 100644 --- a/OpenCL/m00010_a1-optimized.cl +++ b/OpenCL/m00010_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00010_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m00010_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -309,7 +309,7 @@ __kernel void m00010_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -357,9 +357,9 @@ __kernel void m00010_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00020_a1-optimized.cl b/OpenCL/m00020_a1-optimized.cl index 925ae9126..82bfc1fb1 100644 --- a/OpenCL/m00020_a1-optimized.cl +++ b/OpenCL/m00020_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00020_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m00020_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -287,7 +287,7 @@ __kernel void m00020_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -335,9 +335,9 @@ __kernel void m00020_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00030_a1-optimized.cl b/OpenCL/m00030_a1-optimized.cl index 946996754..fdbd26dcb 100644 --- a/OpenCL/m00030_a1-optimized.cl +++ b/OpenCL/m00030_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00030_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m00030_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -314,7 +314,7 @@ __kernel void m00030_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -362,9 +362,9 @@ __kernel void m00030_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00040_a1-optimized.cl b/OpenCL/m00040_a1-optimized.cl index b6e2e07d2..6e544030d 100644 --- a/OpenCL/m00040_a1-optimized.cl +++ b/OpenCL/m00040_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00040_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m00040_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -292,7 +292,7 @@ __kernel void m00040_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -340,9 +340,9 @@ __kernel void m00040_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00050_a1-optimized.cl b/OpenCL/m00050_a1-optimized.cl index 923c9147a..c29f16dae 100644 --- a/OpenCL/m00050_a1-optimized.cl +++ b/OpenCL/m00050_a1-optimized.cl @@ -126,7 +126,7 @@ __kernel void m00050_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -162,9 +162,9 @@ __kernel void m00050_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -299,7 +299,7 @@ __kernel void m00050_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -347,9 +347,9 @@ __kernel void m00050_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00060_a1-optimized.cl b/OpenCL/m00060_a1-optimized.cl index 54752657e..6a5b76819 100644 --- a/OpenCL/m00060_a1-optimized.cl +++ b/OpenCL/m00060_a1-optimized.cl @@ -126,7 +126,7 @@ __kernel void m00060_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -193,9 +193,9 @@ __kernel void m00060_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -304,7 +304,7 @@ __kernel void m00060_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -383,9 +383,9 @@ __kernel void m00060_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00100_a1-optimized.cl b/OpenCL/m00100_a1-optimized.cl index 05d592a23..bc9e9552b 100644 --- a/OpenCL/m00100_a1-optimized.cl +++ b/OpenCL/m00100_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00100_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -48,9 +48,9 @@ __kernel void m00100_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -278,7 +278,7 @@ __kernel void m00100_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -304,9 +304,9 @@ __kernel void m00100_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00110_a1-optimized.cl b/OpenCL/m00110_a1-optimized.cl index 0b975b11f..aba132e5d 100644 --- a/OpenCL/m00110_a1-optimized.cl +++ b/OpenCL/m00110_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00110_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m00110_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -353,7 +353,7 @@ __kernel void m00110_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -407,9 +407,9 @@ __kernel void m00110_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00120_a1-optimized.cl b/OpenCL/m00120_a1-optimized.cl index 3937e11b4..15762357c 100644 --- a/OpenCL/m00120_a1-optimized.cl +++ b/OpenCL/m00120_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00120_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -88,9 +88,9 @@ __kernel void m00120_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -343,7 +343,7 @@ __kernel void m00120_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -397,9 +397,9 @@ __kernel void m00120_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00130_a1-optimized.cl b/OpenCL/m00130_a1-optimized.cl index fe8528750..98ce3989f 100644 --- a/OpenCL/m00130_a1-optimized.cl +++ b/OpenCL/m00130_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00130_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m00130_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -358,7 +358,7 @@ __kernel void m00130_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -412,9 +412,9 @@ __kernel void m00130_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00140_a1-optimized.cl b/OpenCL/m00140_a1-optimized.cl index 2738b4193..fe448f032 100644 --- a/OpenCL/m00140_a1-optimized.cl +++ b/OpenCL/m00140_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00140_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m00140_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -336,7 +336,7 @@ __kernel void m00140_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -390,9 +390,9 @@ __kernel void m00140_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00150_a1-optimized.cl b/OpenCL/m00150_a1-optimized.cl index 289dd77fe..606dfdc1b 100644 --- a/OpenCL/m00150_a1-optimized.cl +++ b/OpenCL/m00150_a1-optimized.cl @@ -130,7 +130,7 @@ __kernel void m00150_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -166,9 +166,9 @@ __kernel void m00150_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -320,7 +320,7 @@ __kernel void m00150_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -368,9 +368,9 @@ __kernel void m00150_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00160_a1-optimized.cl b/OpenCL/m00160_a1-optimized.cl index 71df47c64..0f5025896 100644 --- a/OpenCL/m00160_a1-optimized.cl +++ b/OpenCL/m00160_a1-optimized.cl @@ -130,7 +130,7 @@ __kernel void m00160_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -197,9 +197,9 @@ __kernel void m00160_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -327,7 +327,7 @@ __kernel void m00160_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -406,9 +406,9 @@ __kernel void m00160_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00200_a1-optimized.cl b/OpenCL/m00200_a1-optimized.cl index 296211cb9..ca6800431 100644 --- a/OpenCL/m00200_a1-optimized.cl +++ b/OpenCL/m00200_a1-optimized.cl @@ -41,7 +41,7 @@ __kernel void m00200_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -49,9 +49,9 @@ __kernel void m00200_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -229,7 +229,7 @@ __kernel void m00200_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -249,9 +249,9 @@ __kernel void m00200_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00300_a1-optimized.cl b/OpenCL/m00300_a1-optimized.cl index a961f8df0..1d1becacb 100644 --- a/OpenCL/m00300_a1-optimized.cl +++ b/OpenCL/m00300_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00300_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -48,9 +48,9 @@ __kernel void m00300_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -403,7 +403,7 @@ __kernel void m00300_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -429,9 +429,9 @@ __kernel void m00300_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m00600_a1-optimized.cl b/OpenCL/m00600_a1-optimized.cl index dd0c82b47..be4658f9b 100644 --- a/OpenCL/m00600_a1-optimized.cl +++ b/OpenCL/m00600_a1-optimized.cl @@ -142,7 +142,7 @@ __kernel void m00600_m04 (KERN_ATTR_ESALT (blake2_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; u64 tmp_h[8]; u64 tmp_t[2]; @@ -168,7 +168,7 @@ __kernel void m00600_m04 (KERN_ATTR_ESALT (blake2_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; const u32x out_len = pw_l_len + pw_r_len; @@ -309,7 +309,7 @@ __kernel void m00600_s04 (KERN_ATTR_ESALT (blake2_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; u64 tmp_h[8]; u64 tmp_t[2]; @@ -347,7 +347,7 @@ __kernel void m00600_s04 (KERN_ATTR_ESALT (blake2_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; const u32x out_len = pw_l_len + pw_r_len; diff --git a/OpenCL/m00900_a1-optimized.cl b/OpenCL/m00900_a1-optimized.cl index fd06288f4..8f509af4d 100644 --- a/OpenCL/m00900_a1-optimized.cl +++ b/OpenCL/m00900_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m00900_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -48,9 +48,9 @@ __kernel void m00900_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -215,7 +215,7 @@ __kernel void m00900_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -235,9 +235,9 @@ __kernel void m00900_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01000_a1-optimized.cl b/OpenCL/m01000_a1-optimized.cl index a9d95f55e..1fd29bb6d 100644 --- a/OpenCL/m01000_a1-optimized.cl +++ b/OpenCL/m01000_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m01000_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -48,9 +48,9 @@ __kernel void m01000_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -221,7 +221,7 @@ __kernel void m01000_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -241,9 +241,9 @@ __kernel void m01000_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01100_a1-optimized.cl b/OpenCL/m01100_a1-optimized.cl index 76fcfc4cc..6a2f3bb83 100644 --- a/OpenCL/m01100_a1-optimized.cl +++ b/OpenCL/m01100_a1-optimized.cl @@ -33,7 +33,7 @@ __kernel void m01100_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -70,9 +70,9 @@ __kernel void m01100_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -314,7 +314,7 @@ __kernel void m01100_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -363,9 +363,9 @@ __kernel void m01100_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01300_a1-optimized.cl b/OpenCL/m01300_a1-optimized.cl index b0ca94e2b..b1d627c83 100644 --- a/OpenCL/m01300_a1-optimized.cl +++ b/OpenCL/m01300_a1-optimized.cl @@ -53,7 +53,7 @@ __kernel void m01300_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -61,9 +61,9 @@ __kernel void m01300_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -266,7 +266,7 @@ __kernel void m01300_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -302,9 +302,9 @@ __kernel void m01300_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01400_a1-optimized.cl b/OpenCL/m01400_a1-optimized.cl index b7799853a..b9d104563 100644 --- a/OpenCL/m01400_a1-optimized.cl +++ b/OpenCL/m01400_a1-optimized.cl @@ -54,7 +54,7 @@ __kernel void m01400_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -62,9 +62,9 @@ __kernel void m01400_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -267,7 +267,7 @@ __kernel void m01400_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -305,9 +305,9 @@ __kernel void m01400_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01410_a1-optimized.cl b/OpenCL/m01410_a1-optimized.cl index 6d4a8ba0a..949519ddc 100644 --- a/OpenCL/m01410_a1-optimized.cl +++ b/OpenCL/m01410_a1-optimized.cl @@ -54,7 +54,7 @@ __kernel void m01410_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -90,9 +90,9 @@ __kernel void m01410_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -342,7 +342,7 @@ __kernel void m01410_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -408,9 +408,9 @@ __kernel void m01410_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01420_a1-optimized.cl b/OpenCL/m01420_a1-optimized.cl index fbdc92729..291f3c78f 100644 --- a/OpenCL/m01420_a1-optimized.cl +++ b/OpenCL/m01420_a1-optimized.cl @@ -54,7 +54,7 @@ __kernel void m01420_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -90,9 +90,9 @@ __kernel void m01420_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -320,7 +320,7 @@ __kernel void m01420_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -386,9 +386,9 @@ __kernel void m01420_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01430_a1-optimized.cl b/OpenCL/m01430_a1-optimized.cl index 0b36abb50..24341d471 100644 --- a/OpenCL/m01430_a1-optimized.cl +++ b/OpenCL/m01430_a1-optimized.cl @@ -54,7 +54,7 @@ __kernel void m01430_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -90,9 +90,9 @@ __kernel void m01430_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -347,7 +347,7 @@ __kernel void m01430_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -413,9 +413,9 @@ __kernel void m01430_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01440_a1-optimized.cl b/OpenCL/m01440_a1-optimized.cl index 892ea7acd..24ce60dde 100644 --- a/OpenCL/m01440_a1-optimized.cl +++ b/OpenCL/m01440_a1-optimized.cl @@ -54,7 +54,7 @@ __kernel void m01440_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -90,9 +90,9 @@ __kernel void m01440_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -325,7 +325,7 @@ __kernel void m01440_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -391,9 +391,9 @@ __kernel void m01440_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01450_a1-optimized.cl b/OpenCL/m01450_a1-optimized.cl index 1f9afaf19..e49ee66f1 100644 --- a/OpenCL/m01450_a1-optimized.cl +++ b/OpenCL/m01450_a1-optimized.cl @@ -142,7 +142,7 @@ __kernel void m01450_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -178,9 +178,9 @@ __kernel void m01450_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -332,7 +332,7 @@ __kernel void m01450_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -380,9 +380,9 @@ __kernel void m01450_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01460_a1-optimized.cl b/OpenCL/m01460_a1-optimized.cl index f6ec20939..31519e26b 100644 --- a/OpenCL/m01460_a1-optimized.cl +++ b/OpenCL/m01460_a1-optimized.cl @@ -142,7 +142,7 @@ __kernel void m01460_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -209,9 +209,9 @@ __kernel void m01460_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -339,7 +339,7 @@ __kernel void m01460_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -418,9 +418,9 @@ __kernel void m01460_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01500_a1-pure.cl b/OpenCL/m01500_a1-pure.cl index f184876f6..cba267b9c 100644 --- a/OpenCL/m01500_a1-pure.cl +++ b/OpenCL/m01500_a1-pure.cl @@ -537,7 +537,7 @@ __kernel void m01500_mxx (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -551,9 +551,9 @@ __kernel void m01500_mxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -700,7 +700,7 @@ __kernel void m01500_sxx (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -726,9 +726,9 @@ __kernel void m01500_sxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01700_a1-optimized.cl b/OpenCL/m01700_a1-optimized.cl index db6cbe43b..0892303c6 100644 --- a/OpenCL/m01700_a1-optimized.cl +++ b/OpenCL/m01700_a1-optimized.cl @@ -163,7 +163,7 @@ __kernel void m01700_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -171,9 +171,9 @@ __kernel void m01700_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -322,7 +322,7 @@ __kernel void m01700_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -342,9 +342,9 @@ __kernel void m01700_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01710_a1-optimized.cl b/OpenCL/m01710_a1-optimized.cl index e1ad416be..0d26ad29e 100644 --- a/OpenCL/m01710_a1-optimized.cl +++ b/OpenCL/m01710_a1-optimized.cl @@ -163,7 +163,7 @@ __kernel void m01710_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -199,9 +199,9 @@ __kernel void m01710_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -397,7 +397,7 @@ __kernel void m01710_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -445,9 +445,9 @@ __kernel void m01710_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01720_a1-optimized.cl b/OpenCL/m01720_a1-optimized.cl index 98596c3ff..e98e42d8f 100644 --- a/OpenCL/m01720_a1-optimized.cl +++ b/OpenCL/m01720_a1-optimized.cl @@ -163,7 +163,7 @@ __kernel void m01720_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -199,9 +199,9 @@ __kernel void m01720_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -375,7 +375,7 @@ __kernel void m01720_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -423,9 +423,9 @@ __kernel void m01720_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01730_a1-optimized.cl b/OpenCL/m01730_a1-optimized.cl index 9167e86ba..cf64fa6c1 100644 --- a/OpenCL/m01730_a1-optimized.cl +++ b/OpenCL/m01730_a1-optimized.cl @@ -163,7 +163,7 @@ __kernel void m01730_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -199,9 +199,9 @@ __kernel void m01730_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -402,7 +402,7 @@ __kernel void m01730_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -450,9 +450,9 @@ __kernel void m01730_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01740_a1-optimized.cl b/OpenCL/m01740_a1-optimized.cl index d6e35b0a7..2b9a81e46 100644 --- a/OpenCL/m01740_a1-optimized.cl +++ b/OpenCL/m01740_a1-optimized.cl @@ -163,7 +163,7 @@ __kernel void m01740_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -199,9 +199,9 @@ __kernel void m01740_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -380,7 +380,7 @@ __kernel void m01740_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -428,9 +428,9 @@ __kernel void m01740_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01750_a1-optimized.cl b/OpenCL/m01750_a1-optimized.cl index 465e2eb8f..7dd1a9475 100644 --- a/OpenCL/m01750_a1-optimized.cl +++ b/OpenCL/m01750_a1-optimized.cl @@ -216,7 +216,7 @@ __kernel void m01750_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -252,9 +252,9 @@ __kernel void m01750_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -411,7 +411,7 @@ __kernel void m01750_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -459,9 +459,9 @@ __kernel void m01750_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m01760_a1-optimized.cl b/OpenCL/m01760_a1-optimized.cl index f4b13a4dc..08207c579 100644 --- a/OpenCL/m01760_a1-optimized.cl +++ b/OpenCL/m01760_a1-optimized.cl @@ -216,7 +216,7 @@ __kernel void m01760_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -283,9 +283,9 @@ __kernel void m01760_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -418,7 +418,7 @@ __kernel void m01760_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -497,9 +497,9 @@ __kernel void m01760_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m02400_a1-optimized.cl b/OpenCL/m02400_a1-optimized.cl index c10474699..8c347304a 100644 --- a/OpenCL/m02400_a1-optimized.cl +++ b/OpenCL/m02400_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m02400_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -48,9 +48,9 @@ __kernel void m02400_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -273,7 +273,7 @@ __kernel void m02400_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -293,9 +293,9 @@ __kernel void m02400_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m02410_a1-optimized.cl b/OpenCL/m02410_a1-optimized.cl index ddb12b61b..6f533b689 100644 --- a/OpenCL/m02410_a1-optimized.cl +++ b/OpenCL/m02410_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m02410_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m02410_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -348,7 +348,7 @@ __kernel void m02410_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -396,9 +396,9 @@ __kernel void m02410_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m02610_a1-optimized.cl b/OpenCL/m02610_a1-optimized.cl index 56e63df5a..9bf208976 100644 --- a/OpenCL/m02610_a1-optimized.cl +++ b/OpenCL/m02610_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m02610_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -105,9 +105,9 @@ __kernel void m02610_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -412,7 +412,7 @@ __kernel void m02610_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -460,9 +460,9 @@ __kernel void m02610_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m02710_a1-optimized.cl b/OpenCL/m02710_a1-optimized.cl index 2dc1c521d..1dde3fa29 100644 --- a/OpenCL/m02710_a1-optimized.cl +++ b/OpenCL/m02710_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m02710_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -105,9 +105,9 @@ __kernel void m02710_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -497,7 +497,7 @@ __kernel void m02710_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -545,9 +545,9 @@ __kernel void m02710_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m02810_a1-optimized.cl b/OpenCL/m02810_a1-optimized.cl index 4d74d5c81..230993c22 100644 --- a/OpenCL/m02810_a1-optimized.cl +++ b/OpenCL/m02810_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m02810_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -105,9 +105,9 @@ __kernel void m02810_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -498,7 +498,7 @@ __kernel void m02810_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -546,9 +546,9 @@ __kernel void m02810_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m03000_a1-pure.cl b/OpenCL/m03000_a1-pure.cl index 63a007ac9..80fdc0f33 100644 --- a/OpenCL/m03000_a1-pure.cl +++ b/OpenCL/m03000_a1-pure.cl @@ -547,7 +547,7 @@ __kernel void m03000_mxx (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -555,9 +555,9 @@ __kernel void m03000_mxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -711,7 +711,7 @@ __kernel void m03000_sxx (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -731,9 +731,9 @@ __kernel void m03000_sxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m03100_a1-optimized.cl b/OpenCL/m03100_a1-optimized.cl index e8f427097..546758013 100644 --- a/OpenCL/m03100_a1-optimized.cl +++ b/OpenCL/m03100_a1-optimized.cl @@ -543,7 +543,7 @@ __kernel void m03100_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -569,9 +569,9 @@ __kernel void m03100_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; const u32x salt_word_len = (salt_len + pw_len) * 2; @@ -826,7 +826,7 @@ __kernel void m03100_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -864,9 +864,9 @@ __kernel void m03100_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; const u32x salt_word_len = (salt_len + pw_len) * 2; diff --git a/OpenCL/m03710_a1-optimized.cl b/OpenCL/m03710_a1-optimized.cl index 9d39daba9..afef29fdb 100644 --- a/OpenCL/m03710_a1-optimized.cl +++ b/OpenCL/m03710_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m03710_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -107,9 +107,9 @@ __kernel void m03710_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -443,7 +443,7 @@ __kernel void m03710_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -493,9 +493,9 @@ __kernel void m03710_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m03800_a1-optimized.cl b/OpenCL/m03800_a1-optimized.cl index 115ab79af..aa364d137 100644 --- a/OpenCL/m03800_a1-optimized.cl +++ b/OpenCL/m03800_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m03800_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m03800_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -339,7 +339,7 @@ __kernel void m03800_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -387,9 +387,9 @@ __kernel void m03800_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m03910_a1-optimized.cl b/OpenCL/m03910_a1-optimized.cl index 82f319050..b5acc65cf 100644 --- a/OpenCL/m03910_a1-optimized.cl +++ b/OpenCL/m03910_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m03910_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -105,9 +105,9 @@ __kernel void m03910_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -498,7 +498,7 @@ __kernel void m03910_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -546,9 +546,9 @@ __kernel void m03910_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m04010_a1-optimized.cl b/OpenCL/m04010_a1-optimized.cl index 32475ea12..c0bec22b4 100644 --- a/OpenCL/m04010_a1-optimized.cl +++ b/OpenCL/m04010_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m04010_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -107,9 +107,9 @@ __kernel void m04010_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -468,7 +468,7 @@ __kernel void m04010_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -518,9 +518,9 @@ __kernel void m04010_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m04110_a1-optimized.cl b/OpenCL/m04110_a1-optimized.cl index 6120259dc..6ef8064f8 100644 --- a/OpenCL/m04110_a1-optimized.cl +++ b/OpenCL/m04110_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m04110_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -132,9 +132,9 @@ __kernel void m04110_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -515,7 +515,7 @@ __kernel void m04110_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -590,9 +590,9 @@ __kernel void m04110_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m04310_a1-optimized.cl b/OpenCL/m04310_a1-optimized.cl index 76673020d..b27745893 100644 --- a/OpenCL/m04310_a1-optimized.cl +++ b/OpenCL/m04310_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m04310_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -105,9 +105,9 @@ __kernel void m04310_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -412,7 +412,7 @@ __kernel void m04310_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -460,9 +460,9 @@ __kernel void m04310_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m04400_a1-optimized.cl b/OpenCL/m04400_a1-optimized.cl index accebb9eb..ae4b1324e 100644 --- a/OpenCL/m04400_a1-optimized.cl +++ b/OpenCL/m04400_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m04400_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -77,9 +77,9 @@ __kernel void m04400_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -437,7 +437,7 @@ __kernel void m04400_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -457,9 +457,9 @@ __kernel void m04400_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m04500_a1-optimized.cl b/OpenCL/m04500_a1-optimized.cl index a447ef3f8..e4b6fa126 100644 --- a/OpenCL/m04500_a1-optimized.cl +++ b/OpenCL/m04500_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m04500_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -77,9 +77,9 @@ __kernel void m04500_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -464,7 +464,7 @@ __kernel void m04500_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -490,9 +490,9 @@ __kernel void m04500_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m04520_a1-optimized.cl b/OpenCL/m04520_a1-optimized.cl index cc4f6d6da..b967b9f3b 100644 --- a/OpenCL/m04520_a1-optimized.cl +++ b/OpenCL/m04520_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m04520_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -105,9 +105,9 @@ __kernel void m04520_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -695,7 +695,7 @@ __kernel void m04520_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -743,9 +743,9 @@ __kernel void m04520_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m04700_a1-optimized.cl b/OpenCL/m04700_a1-optimized.cl index 928938814..747de70e6 100644 --- a/OpenCL/m04700_a1-optimized.cl +++ b/OpenCL/m04700_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m04700_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -77,9 +77,9 @@ __kernel void m04700_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -417,7 +417,7 @@ __kernel void m04700_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -443,9 +443,9 @@ __kernel void m04700_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m04800_a1-optimized.cl b/OpenCL/m04800_a1-optimized.cl index 7541133a9..679b28dc1 100644 --- a/OpenCL/m04800_a1-optimized.cl +++ b/OpenCL/m04800_a1-optimized.cl @@ -51,7 +51,7 @@ __kernel void m04800_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -73,9 +73,9 @@ __kernel void m04800_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; const u32x pw_salt_len = pw_len + salt_len; @@ -316,7 +316,7 @@ __kernel void m04800_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -350,9 +350,9 @@ __kernel void m04800_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; const u32x pw_salt_len = pw_len + salt_len; diff --git a/OpenCL/m04900_a1-optimized.cl b/OpenCL/m04900_a1-optimized.cl index 1b286a00d..fcbd1bbc0 100644 --- a/OpenCL/m04900_a1-optimized.cl +++ b/OpenCL/m04900_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m04900_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m04900_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -380,7 +380,7 @@ __kernel void m04900_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -434,9 +434,9 @@ __kernel void m04900_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m05100_a1-optimized.cl b/OpenCL/m05100_a1-optimized.cl index 09c9ac35c..80bed0f23 100644 --- a/OpenCL/m05100_a1-optimized.cl +++ b/OpenCL/m05100_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m05100_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -48,9 +48,9 @@ __kernel void m05100_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -243,7 +243,7 @@ __kernel void m05100_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -263,9 +263,9 @@ __kernel void m05100_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m05300_a1-optimized.cl b/OpenCL/m05300_a1-optimized.cl index fc1ff9472..1738a951a 100644 --- a/OpenCL/m05300_a1-optimized.cl +++ b/OpenCL/m05300_a1-optimized.cl @@ -146,7 +146,7 @@ __kernel void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -161,9 +161,9 @@ __kernel void m05300_m04 (KERN_ATTR_ESALT (ikepsk_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -381,7 +381,7 @@ __kernel void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -408,9 +408,9 @@ __kernel void m05300_s04 (KERN_ATTR_ESALT (ikepsk_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m05400_a1-optimized.cl b/OpenCL/m05400_a1-optimized.cl index eff432407..9d2bcf127 100644 --- a/OpenCL/m05400_a1-optimized.cl +++ b/OpenCL/m05400_a1-optimized.cl @@ -150,7 +150,7 @@ __kernel void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -165,9 +165,9 @@ __kernel void m05400_m04 (KERN_ATTR_ESALT (ikepsk_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -402,7 +402,7 @@ __kernel void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -429,9 +429,9 @@ __kernel void m05400_s04 (KERN_ATTR_ESALT (ikepsk_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m05500_a1-optimized.cl b/OpenCL/m05500_a1-optimized.cl index 9420df0fe..6ae857bf8 100644 --- a/OpenCL/m05500_a1-optimized.cl +++ b/OpenCL/m05500_a1-optimized.cl @@ -544,7 +544,7 @@ __kernel void m05500_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -560,9 +560,9 @@ __kernel void m05500_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -810,7 +810,7 @@ __kernel void m05500_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -838,9 +838,9 @@ __kernel void m05500_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m05600_a1-optimized.cl b/OpenCL/m05600_a1-optimized.cl index de20b8632..8aeb5fb8e 100644 --- a/OpenCL/m05600_a1-optimized.cl +++ b/OpenCL/m05600_a1-optimized.cl @@ -153,7 +153,7 @@ __kernel void m05600_m04 (KERN_ATTR_ESALT (netntlm_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -161,9 +161,9 @@ __kernel void m05600_m04 (KERN_ATTR_ESALT (netntlm_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -450,7 +450,7 @@ __kernel void m05600_s04 (KERN_ATTR_ESALT (netntlm_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -470,9 +470,9 @@ __kernel void m05600_s04 (KERN_ATTR_ESALT (netntlm_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m06000_a1-optimized.cl b/OpenCL/m06000_a1-optimized.cl index 92c52b1bc..cd6e25885 100644 --- a/OpenCL/m06000_a1-optimized.cl +++ b/OpenCL/m06000_a1-optimized.cl @@ -46,7 +46,7 @@ __kernel void m06000_m04 (KERN_ATTR_ESALT (netntlm_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -54,9 +54,9 @@ __kernel void m06000_m04 (KERN_ATTR_ESALT (netntlm_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -194,7 +194,7 @@ __kernel void m06000_s04 (KERN_ATTR_ESALT (netntlm_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -214,9 +214,9 @@ __kernel void m06000_s04 (KERN_ATTR_ESALT (netntlm_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m06100_a1-optimized.cl b/OpenCL/m06100_a1-optimized.cl index 005d7f9ea..abaa6b039 100644 --- a/OpenCL/m06100_a1-optimized.cl +++ b/OpenCL/m06100_a1-optimized.cl @@ -76,7 +76,7 @@ __kernel void m06100_m04 (KERN_ATTR_ESALT (netntlm_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -84,9 +84,9 @@ __kernel void m06100_m04 (KERN_ATTR_ESALT (netntlm_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -265,7 +265,7 @@ __kernel void m06100_s04 (KERN_ATTR_ESALT (netntlm_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -285,9 +285,9 @@ __kernel void m06100_s04 (KERN_ATTR_ESALT (netntlm_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m06900_a1-optimized.cl b/OpenCL/m06900_a1-optimized.cl index 2677a7ad5..ef77bcdcf 100644 --- a/OpenCL/m06900_a1-optimized.cl +++ b/OpenCL/m06900_a1-optimized.cl @@ -739,7 +739,7 @@ __kernel void m06900_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -747,9 +747,9 @@ __kernel void m06900_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -1009,7 +1009,7 @@ __kernel void m06900_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -1029,9 +1029,9 @@ __kernel void m06900_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m07000_a1-optimized.cl b/OpenCL/m07000_a1-optimized.cl index e61337bc0..da9ee569e 100644 --- a/OpenCL/m07000_a1-optimized.cl +++ b/OpenCL/m07000_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m07000_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -62,9 +62,9 @@ __kernel void m07000_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -362,7 +362,7 @@ __kernel void m07000_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -402,9 +402,9 @@ __kernel void m07000_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m07300_a1-optimized.cl b/OpenCL/m07300_a1-optimized.cl index 2d3ac2150..6f096a277 100644 --- a/OpenCL/m07300_a1-optimized.cl +++ b/OpenCL/m07300_a1-optimized.cl @@ -130,7 +130,7 @@ __kernel void m07300_m04 (KERN_ATTR_ESALT (rakp_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -144,9 +144,9 @@ __kernel void m07300_m04 (KERN_ATTR_ESALT (rakp_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -325,7 +325,7 @@ __kernel void m07300_s04 (KERN_ATTR_ESALT (rakp_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -351,9 +351,9 @@ __kernel void m07300_s04 (KERN_ATTR_ESALT (rakp_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m07500_a1-optimized.cl b/OpenCL/m07500_a1-optimized.cl index 366b714be..e732afb77 100644 --- a/OpenCL/m07500_a1-optimized.cl +++ b/OpenCL/m07500_a1-optimized.cl @@ -410,7 +410,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -448,9 +448,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -566,7 +566,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -604,9 +604,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m07700_a1-optimized.cl b/OpenCL/m07700_a1-optimized.cl index 7f2dc993a..f0d4efbd4 100644 --- a/OpenCL/m07700_a1-optimized.cl +++ b/OpenCL/m07700_a1-optimized.cl @@ -173,7 +173,7 @@ __kernel void m07700_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -198,9 +198,9 @@ __kernel void m07700_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -380,7 +380,7 @@ __kernel void m07700_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -417,9 +417,9 @@ __kernel void m07700_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m07701_a1-optimized.cl b/OpenCL/m07701_a1-optimized.cl index f275beff4..4f4e909a1 100644 --- a/OpenCL/m07701_a1-optimized.cl +++ b/OpenCL/m07701_a1-optimized.cl @@ -173,7 +173,7 @@ __kernel void m07701_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -198,9 +198,9 @@ __kernel void m07701_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -380,7 +380,7 @@ __kernel void m07701_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -417,9 +417,9 @@ __kernel void m07701_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m07800_a1-optimized.cl b/OpenCL/m07800_a1-optimized.cl index 2d394c3e6..b9cbd893c 100644 --- a/OpenCL/m07800_a1-optimized.cl +++ b/OpenCL/m07800_a1-optimized.cl @@ -79,7 +79,7 @@ __kernel void m07800_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -104,9 +104,9 @@ __kernel void m07800_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -387,7 +387,7 @@ __kernel void m07800_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -424,9 +424,9 @@ __kernel void m07800_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m07801_a1-optimized.cl b/OpenCL/m07801_a1-optimized.cl index 878919857..37b29634e 100644 --- a/OpenCL/m07801_a1-optimized.cl +++ b/OpenCL/m07801_a1-optimized.cl @@ -79,7 +79,7 @@ __kernel void m07801_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -104,9 +104,9 @@ __kernel void m07801_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -387,7 +387,7 @@ __kernel void m07801_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -424,9 +424,9 @@ __kernel void m07801_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m08000_a1-optimized.cl b/OpenCL/m08000_a1-optimized.cl index 55caae104..0b1d66424 100644 --- a/OpenCL/m08000_a1-optimized.cl +++ b/OpenCL/m08000_a1-optimized.cl @@ -306,7 +306,7 @@ __kernel void m08000_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -314,9 +314,9 @@ __kernel void m08000_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -531,7 +531,7 @@ __kernel void m08000_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -551,9 +551,9 @@ __kernel void m08000_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m08100_a1-optimized.cl b/OpenCL/m08100_a1-optimized.cl index 2c58a4a42..5366b9fa8 100644 --- a/OpenCL/m08100_a1-optimized.cl +++ b/OpenCL/m08100_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m08100_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -71,9 +71,9 @@ __kernel void m08100_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -323,7 +323,7 @@ __kernel void m08100_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -360,9 +360,9 @@ __kernel void m08100_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m08300_a1-optimized.cl b/OpenCL/m08300_a1-optimized.cl index 0e3e1dc60..0d330e27e 100644 --- a/OpenCL/m08300_a1-optimized.cl +++ b/OpenCL/m08300_a1-optimized.cl @@ -41,7 +41,7 @@ __kernel void m08300_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -83,9 +83,9 @@ __kernel void m08300_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -358,7 +358,7 @@ __kernel void m08300_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -412,9 +412,9 @@ __kernel void m08300_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m08400_a1-optimized.cl b/OpenCL/m08400_a1-optimized.cl index 89fc0d66f..63419d6b1 100644 --- a/OpenCL/m08400_a1-optimized.cl +++ b/OpenCL/m08400_a1-optimized.cl @@ -70,7 +70,7 @@ __kernel void m08400_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -101,9 +101,9 @@ __kernel void m08400_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -382,7 +382,7 @@ __kernel void m08400_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -425,9 +425,9 @@ __kernel void m08400_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m08500_a1-pure.cl b/OpenCL/m08500_a1-pure.cl index 31efadece..ec3ee5c03 100644 --- a/OpenCL/m08500_a1-pure.cl +++ b/OpenCL/m08500_a1-pure.cl @@ -577,7 +577,7 @@ __kernel void m08500_mxx (KERN_ATTR_BASIC ()) pw_buf1[2] = 0; pw_buf1[3] = 0; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -594,9 +594,9 @@ __kernel void m08500_mxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -733,7 +733,7 @@ __kernel void m08500_sxx (KERN_ATTR_BASIC ()) pw_buf1[2] = 0; pw_buf1[3] = 0; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -762,9 +762,9 @@ __kernel void m08500_sxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m08600_a1-pure.cl b/OpenCL/m08600_a1-pure.cl index 5be8800fd..9ad4d7bf1 100644 --- a/OpenCL/m08600_a1-pure.cl +++ b/OpenCL/m08600_a1-pure.cl @@ -268,7 +268,7 @@ __kernel void m08600_mxx (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -276,9 +276,9 @@ __kernel void m08600_mxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -392,7 +392,7 @@ __kernel void m08600_sxx (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -412,9 +412,9 @@ __kernel void m08600_sxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m08700_a1-optimized.cl b/OpenCL/m08700_a1-optimized.cl index dffb07668..954b52ff1 100644 --- a/OpenCL/m08700_a1-optimized.cl +++ b/OpenCL/m08700_a1-optimized.cl @@ -316,7 +316,7 @@ __kernel void m08700_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -331,9 +331,9 @@ __kernel void m08700_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -563,7 +563,7 @@ __kernel void m08700_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -590,9 +590,9 @@ __kernel void m08700_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m09700_a1-optimized.cl b/OpenCL/m09700_a1-optimized.cl index 0dfce6421..2fedbafdc 100644 --- a/OpenCL/m09700_a1-optimized.cl +++ b/OpenCL/m09700_a1-optimized.cl @@ -509,7 +509,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -549,9 +549,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -756,7 +756,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -808,9 +808,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09700_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m09710_a1-optimized.cl b/OpenCL/m09710_a1-optimized.cl index 98dc83198..5cd5a923d 100644 --- a/OpenCL/m09710_a1-optimized.cl +++ b/OpenCL/m09710_a1-optimized.cl @@ -165,7 +165,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -194,9 +194,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -351,7 +351,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -392,9 +392,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09710_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m09720_a1-optimized.cl b/OpenCL/m09720_a1-optimized.cl index baf93992c..956d9b5a6 100644 --- a/OpenCL/m09720_a1-optimized.cl +++ b/OpenCL/m09720_a1-optimized.cl @@ -385,7 +385,7 @@ __kernel void m09720_m04 (KERN_ATTR_ESALT (oldoffice01_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -404,9 +404,9 @@ __kernel void m09720_m04 (KERN_ATTR_ESALT (oldoffice01_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -541,7 +541,7 @@ __kernel void m09720_s04 (KERN_ATTR_ESALT (oldoffice01_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -572,9 +572,9 @@ __kernel void m09720_s04 (KERN_ATTR_ESALT (oldoffice01_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m09800_a1-optimized.cl b/OpenCL/m09800_a1-optimized.cl index 006895437..ba64ac196 100644 --- a/OpenCL/m09800_a1-optimized.cl +++ b/OpenCL/m09800_a1-optimized.cl @@ -165,7 +165,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -205,9 +205,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -415,7 +415,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -467,9 +467,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09800_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m09810_a1-optimized.cl b/OpenCL/m09810_a1-optimized.cl index 8fb1746d0..e0cce2e77 100644 --- a/OpenCL/m09810_a1-optimized.cl +++ b/OpenCL/m09810_a1-optimized.cl @@ -165,7 +165,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -194,9 +194,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -338,7 +338,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -379,9 +379,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m09810_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m09820_a1-optimized.cl b/OpenCL/m09820_a1-optimized.cl index e37e4f6fc..cc07760a4 100644 --- a/OpenCL/m09820_a1-optimized.cl +++ b/OpenCL/m09820_a1-optimized.cl @@ -41,7 +41,7 @@ __kernel void m09820_m04 (KERN_ATTR_ESALT (oldoffice34_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -60,9 +60,9 @@ __kernel void m09820_m04 (KERN_ATTR_ESALT (oldoffice34_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -225,7 +225,7 @@ __kernel void m09820_s04 (KERN_ATTR_ESALT (oldoffice34_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -256,9 +256,9 @@ __kernel void m09820_s04 (KERN_ATTR_ESALT (oldoffice34_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m09900_a1-optimized.cl b/OpenCL/m09900_a1-optimized.cl index 0f8809199..1d2396f57 100644 --- a/OpenCL/m09900_a1-optimized.cl +++ b/OpenCL/m09900_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m09900_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -48,9 +48,9 @@ __kernel void m09900_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -334,7 +334,7 @@ __kernel void m09900_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -354,9 +354,9 @@ __kernel void m09900_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m10100_a1-optimized.cl b/OpenCL/m10100_a1-optimized.cl index b2ba3cda4..49de592cf 100644 --- a/OpenCL/m10100_a1-optimized.cl +++ b/OpenCL/m10100_a1-optimized.cl @@ -57,7 +57,7 @@ __kernel void m10100_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -79,9 +79,9 @@ __kernel void m10100_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -244,7 +244,7 @@ __kernel void m10100_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -278,9 +278,9 @@ __kernel void m10100_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m10400_a1-optimized.cl b/OpenCL/m10400_a1-optimized.cl index 80002a86d..2b843c695 100644 --- a/OpenCL/m10400_a1-optimized.cl +++ b/OpenCL/m10400_a1-optimized.cl @@ -162,7 +162,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -202,9 +202,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -398,7 +398,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -450,9 +450,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10400_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m10410_a1-optimized.cl b/OpenCL/m10410_a1-optimized.cl index 4b371a541..a5d5a6ed7 100644 --- a/OpenCL/m10410_a1-optimized.cl +++ b/OpenCL/m10410_a1-optimized.cl @@ -162,7 +162,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -202,9 +202,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -304,7 +304,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -356,9 +356,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m10410_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m10420_a1-optimized.cl b/OpenCL/m10420_a1-optimized.cl index c53dd2334..853135537 100644 --- a/OpenCL/m10420_a1-optimized.cl +++ b/OpenCL/m10420_a1-optimized.cl @@ -53,7 +53,7 @@ __kernel void m10420_m04 (KERN_ATTR_ESALT (pdf_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * U_buf @@ -85,9 +85,9 @@ __kernel void m10420_m04 (KERN_ATTR_ESALT (pdf_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -274,7 +274,7 @@ __kernel void m10420_s04 (KERN_ATTR_ESALT (pdf_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * U_buf @@ -318,9 +318,9 @@ __kernel void m10420_s04 (KERN_ATTR_ESALT (pdf_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m10800_a1-optimized.cl b/OpenCL/m10800_a1-optimized.cl index 2596c4ba2..032900925 100644 --- a/OpenCL/m10800_a1-optimized.cl +++ b/OpenCL/m10800_a1-optimized.cl @@ -163,7 +163,7 @@ __kernel void m10800_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -171,9 +171,9 @@ __kernel void m10800_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -322,7 +322,7 @@ __kernel void m10800_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -342,9 +342,9 @@ __kernel void m10800_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m11000_a1-optimized.cl b/OpenCL/m11000_a1-optimized.cl index 54c11ff47..31ede04b1 100644 --- a/OpenCL/m11000_a1-optimized.cl +++ b/OpenCL/m11000_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m11000_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -76,9 +76,9 @@ __kernel void m11000_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -395,7 +395,7 @@ __kernel void m11000_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -443,9 +443,9 @@ __kernel void m11000_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m11100_a1-optimized.cl b/OpenCL/m11100_a1-optimized.cl index 307ae7fb0..4bf7af29c 100644 --- a/OpenCL/m11100_a1-optimized.cl +++ b/OpenCL/m11100_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m11100_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * challenge @@ -103,9 +103,9 @@ __kernel void m11100_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -463,7 +463,7 @@ __kernel void m11100_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * challenge @@ -509,9 +509,9 @@ __kernel void m11100_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m11200_a1-optimized.cl b/OpenCL/m11200_a1-optimized.cl index 33b81e7ce..fccbbb142 100644 --- a/OpenCL/m11200_a1-optimized.cl +++ b/OpenCL/m11200_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m11200_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -60,9 +60,9 @@ __kernel void m11200_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -560,7 +560,7 @@ __kernel void m11200_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -592,9 +592,9 @@ __kernel void m11200_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m11500_a1-optimized.cl b/OpenCL/m11500_a1-optimized.cl index 0c30b1863..d74a0e580 100644 --- a/OpenCL/m11500_a1-optimized.cl +++ b/OpenCL/m11500_a1-optimized.cl @@ -152,7 +152,7 @@ __kernel void m11500_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -166,9 +166,9 @@ __kernel void m11500_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -300,7 +300,7 @@ __kernel void m11500_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -326,9 +326,9 @@ __kernel void m11500_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m11700_a1-optimized.cl b/OpenCL/m11700_a1-optimized.cl index e90c93200..081626ab2 100644 --- a/OpenCL/m11700_a1-optimized.cl +++ b/OpenCL/m11700_a1-optimized.cl @@ -2347,7 +2347,7 @@ __kernel void m11700_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -2355,9 +2355,9 @@ __kernel void m11700_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -2563,7 +2563,7 @@ __kernel void m11700_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -2583,9 +2583,9 @@ __kernel void m11700_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m11800_a1-optimized.cl b/OpenCL/m11800_a1-optimized.cl index 3b419746e..c38c31c82 100644 --- a/OpenCL/m11800_a1-optimized.cl +++ b/OpenCL/m11800_a1-optimized.cl @@ -2347,7 +2347,7 @@ __kernel void m11800_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -2355,9 +2355,9 @@ __kernel void m11800_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -2563,7 +2563,7 @@ __kernel void m11800_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -2583,9 +2583,9 @@ __kernel void m11800_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m12600_a1-optimized.cl b/OpenCL/m12600_a1-optimized.cl index f17d50ddd..6d02b6259 100644 --- a/OpenCL/m12600_a1-optimized.cl +++ b/OpenCL/m12600_a1-optimized.cl @@ -69,7 +69,7 @@ __kernel void m12600_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -92,9 +92,9 @@ __kernel void m12600_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -467,7 +467,7 @@ __kernel void m12600_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -502,9 +502,9 @@ __kernel void m12600_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m13100_a1-optimized.cl b/OpenCL/m13100_a1-optimized.cl index 5b94b059a..0e257946b 100644 --- a/OpenCL/m13100_a1-optimized.cl +++ b/OpenCL/m13100_a1-optimized.cl @@ -591,7 +591,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -618,9 +618,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -738,7 +738,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -765,9 +765,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m13300_a1-optimized.cl b/OpenCL/m13300_a1-optimized.cl index b2bb0e119..d14a4a452 100644 --- a/OpenCL/m13300_a1-optimized.cl +++ b/OpenCL/m13300_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m13300_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -48,9 +48,9 @@ __kernel void m13300_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -285,7 +285,7 @@ __kernel void m13300_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -305,9 +305,9 @@ __kernel void m13300_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m13500_a1-optimized.cl b/OpenCL/m13500_a1-optimized.cl index b8f188bb1..453392b4e 100644 --- a/OpenCL/m13500_a1-optimized.cl +++ b/OpenCL/m13500_a1-optimized.cl @@ -41,7 +41,7 @@ __kernel void m13500_m04 (KERN_ATTR_ESALT (pstoken_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -89,9 +89,9 @@ __kernel void m13500_m04 (KERN_ATTR_ESALT (pstoken_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -518,7 +518,7 @@ __kernel void m13500_s04 (KERN_ATTR_ESALT (pstoken_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -578,9 +578,9 @@ __kernel void m13500_s04 (KERN_ATTR_ESALT (pstoken_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m13800_a1-optimized.cl b/OpenCL/m13800_a1-optimized.cl index 36bb42f94..99a15bf70 100644 --- a/OpenCL/m13800_a1-optimized.cl +++ b/OpenCL/m13800_a1-optimized.cl @@ -420,7 +420,7 @@ __kernel void m13800_m04 (KERN_ATTR_ESALT (win8phone_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -443,9 +443,9 @@ __kernel void m13800_m04 (KERN_ATTR_ESALT (win8phone_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -672,7 +672,7 @@ __kernel void m13800_s04 (KERN_ATTR_ESALT (win8phone_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -707,9 +707,9 @@ __kernel void m13800_s04 (KERN_ATTR_ESALT (win8phone_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m13900_a1-optimized.cl b/OpenCL/m13900_a1-optimized.cl index 5a4960d74..224c7776f 100644 --- a/OpenCL/m13900_a1-optimized.cl +++ b/OpenCL/m13900_a1-optimized.cl @@ -70,7 +70,7 @@ __kernel void m13900_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -90,9 +90,9 @@ __kernel void m13900_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -355,7 +355,7 @@ __kernel void m13900_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -387,9 +387,9 @@ __kernel void m13900_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m14000_a1-pure.cl b/OpenCL/m14000_a1-pure.cl index b15e46878..0642abd86 100644 --- a/OpenCL/m14000_a1-pure.cl +++ b/OpenCL/m14000_a1-pure.cl @@ -544,7 +544,7 @@ __kernel void m14000_mxx (KERN_ATTR_BASIC ()) pw_buf1[2] = 0; pw_buf1[3] = 0; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -561,9 +561,9 @@ __kernel void m14000_mxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -692,7 +692,7 @@ __kernel void m14000_sxx (KERN_ATTR_BASIC ()) pw_buf1[2] = 0; pw_buf1[3] = 0; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -721,9 +721,9 @@ __kernel void m14000_sxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m14100_a1-pure.cl b/OpenCL/m14100_a1-pure.cl index 0c58bd257..07f50acbb 100644 --- a/OpenCL/m14100_a1-pure.cl +++ b/OpenCL/m14100_a1-pure.cl @@ -588,7 +588,7 @@ __kernel void m14100_mxx (KERN_ATTR_BASIC ()) pw_buf1[2] = 0; pw_buf1[3] = 0; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -605,9 +605,9 @@ __kernel void m14100_mxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -771,7 +771,7 @@ __kernel void m14100_sxx (KERN_ATTR_BASIC ()) pw_buf1[2] = 0; pw_buf1[3] = 0; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -800,9 +800,9 @@ __kernel void m14100_sxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m14400_a1-optimized.cl b/OpenCL/m14400_a1-optimized.cl index 3a91768a7..8e92bb064 100644 --- a/OpenCL/m14400_a1-optimized.cl +++ b/OpenCL/m14400_a1-optimized.cl @@ -162,7 +162,7 @@ __kernel void m14400_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -224,9 +224,9 @@ __kernel void m14400_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -497,7 +497,7 @@ __kernel void m14400_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -571,9 +571,9 @@ __kernel void m14400_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m14900_a1-optimized.cl b/OpenCL/m14900_a1-optimized.cl index 16818b0fa..a0630d89e 100644 --- a/OpenCL/m14900_a1-optimized.cl +++ b/OpenCL/m14900_a1-optimized.cl @@ -143,7 +143,7 @@ __kernel void m14900_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -157,9 +157,9 @@ __kernel void m14900_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -298,7 +298,7 @@ __kernel void m14900_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -324,9 +324,9 @@ __kernel void m14900_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m15000_a1-optimized.cl b/OpenCL/m15000_a1-optimized.cl index a0f4f70d4..636c5ae56 100644 --- a/OpenCL/m15000_a1-optimized.cl +++ b/OpenCL/m15000_a1-optimized.cl @@ -163,7 +163,7 @@ __kernel void m15000_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -199,9 +199,9 @@ __kernel void m15000_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -428,7 +428,7 @@ __kernel void m15000_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -476,9 +476,9 @@ __kernel void m15000_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m15400_a1-optimized.cl b/OpenCL/m15400_a1-optimized.cl index d2c0e85d6..13378b279 100644 --- a/OpenCL/m15400_a1-optimized.cl +++ b/OpenCL/m15400_a1-optimized.cl @@ -219,7 +219,7 @@ __kernel void m15400_m04 (KERN_ATTR_ESALT (chacha20_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * Salt prep @@ -247,7 +247,7 @@ __kernel void m15400_m04 (KERN_ATTR_ESALT (chacha20_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; const u32x out_len = pw_l_len + pw_r_len; @@ -349,7 +349,7 @@ __kernel void m15400_s04 (KERN_ATTR_ESALT (chacha20_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * Salt prep @@ -389,7 +389,7 @@ __kernel void m15400_s04 (KERN_ATTR_ESALT (chacha20_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; const u32x out_len = pw_l_len + pw_r_len; diff --git a/OpenCL/m15500_a1-optimized.cl b/OpenCL/m15500_a1-optimized.cl index 476c31409..b85ca9468 100644 --- a/OpenCL/m15500_a1-optimized.cl +++ b/OpenCL/m15500_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m15500_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -62,9 +62,9 @@ __kernel void m15500_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -356,7 +356,7 @@ __kernel void m15500_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -390,9 +390,9 @@ __kernel void m15500_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m16000_a1-pure.cl b/OpenCL/m16000_a1-pure.cl index f267a3ecd..e9b0d284b 100644 --- a/OpenCL/m16000_a1-pure.cl +++ b/OpenCL/m16000_a1-pure.cl @@ -556,7 +556,7 @@ __kernel void m16000_mxx (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -564,9 +564,9 @@ __kernel void m16000_mxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -728,7 +728,7 @@ __kernel void m16000_sxx (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -748,9 +748,9 @@ __kernel void m16000_sxx (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32 pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32 pw_len = pw_l_len + pw_r_len; + const u32 pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m16100_a1-optimized.cl b/OpenCL/m16100_a1-optimized.cl index d44b66884..8ae70c8ff 100644 --- a/OpenCL/m16100_a1-optimized.cl +++ b/OpenCL/m16100_a1-optimized.cl @@ -40,7 +40,7 @@ __kernel void m16100_m04 (KERN_ATTR_ESALT (tacacs_plus_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -66,9 +66,9 @@ __kernel void m16100_m04 (KERN_ATTR_ESALT (tacacs_plus_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -386,7 +386,7 @@ __kernel void m16100_s04 (KERN_ATTR_ESALT (tacacs_plus_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * salt @@ -412,9 +412,9 @@ __kernel void m16100_s04 (KERN_ATTR_ESALT (tacacs_plus_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m16400_a1-optimized.cl b/OpenCL/m16400_a1-optimized.cl index 807ef533c..83b4d2907 100644 --- a/OpenCL/m16400_a1-optimized.cl +++ b/OpenCL/m16400_a1-optimized.cl @@ -42,7 +42,7 @@ __kernel void m16400_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -50,9 +50,9 @@ __kernel void m16400_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -236,7 +236,7 @@ __kernel void m16400_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -256,9 +256,9 @@ __kernel void m16400_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m16600_a1-optimized.cl b/OpenCL/m16600_a1-optimized.cl index dd7502934..512e7de95 100644 --- a/OpenCL/m16600_a1-optimized.cl +++ b/OpenCL/m16600_a1-optimized.cl @@ -92,7 +92,7 @@ __kernel void m16600_m04 (KERN_ATTR_ESALT (electrum_wallet_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -100,9 +100,9 @@ __kernel void m16600_m04 (KERN_ATTR_ESALT (electrum_wallet_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -519,7 +519,7 @@ __kernel void m16600_s04 (KERN_ATTR_ESALT (electrum_wallet_t)) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -527,9 +527,9 @@ __kernel void m16600_s04 (KERN_ATTR_ESALT (electrum_wallet_t)) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m17300_a1-optimized.cl b/OpenCL/m17300_a1-optimized.cl index b77921be2..7b5153c7c 100644 --- a/OpenCL/m17300_a1-optimized.cl +++ b/OpenCL/m17300_a1-optimized.cl @@ -56,7 +56,7 @@ __kernel void m17300_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -64,9 +64,9 @@ __kernel void m17300_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -333,7 +333,7 @@ __kernel void m17300_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -353,9 +353,9 @@ __kernel void m17300_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m17400_a1-optimized.cl b/OpenCL/m17400_a1-optimized.cl index c15386c92..2822073e4 100644 --- a/OpenCL/m17400_a1-optimized.cl +++ b/OpenCL/m17400_a1-optimized.cl @@ -56,7 +56,7 @@ __kernel void m17400_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -64,9 +64,9 @@ __kernel void m17400_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -333,7 +333,7 @@ __kernel void m17400_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -353,9 +353,9 @@ __kernel void m17400_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m17500_a1-optimized.cl b/OpenCL/m17500_a1-optimized.cl index 858d4bc7c..151a0f306 100644 --- a/OpenCL/m17500_a1-optimized.cl +++ b/OpenCL/m17500_a1-optimized.cl @@ -56,7 +56,7 @@ __kernel void m17500_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -64,9 +64,9 @@ __kernel void m17500_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -333,7 +333,7 @@ __kernel void m17500_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -353,9 +353,9 @@ __kernel void m17500_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m17600_a1-optimized.cl b/OpenCL/m17600_a1-optimized.cl index 1da61ddd9..210ce24bd 100644 --- a/OpenCL/m17600_a1-optimized.cl +++ b/OpenCL/m17600_a1-optimized.cl @@ -56,7 +56,7 @@ __kernel void m17600_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -64,9 +64,9 @@ __kernel void m17600_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -333,7 +333,7 @@ __kernel void m17600_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -353,9 +353,9 @@ __kernel void m17600_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m17700_a1-optimized.cl b/OpenCL/m17700_a1-optimized.cl index f0bd3cb8e..e08393d9f 100644 --- a/OpenCL/m17700_a1-optimized.cl +++ b/OpenCL/m17700_a1-optimized.cl @@ -56,7 +56,7 @@ __kernel void m17700_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -64,9 +64,9 @@ __kernel void m17700_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -333,7 +333,7 @@ __kernel void m17700_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -353,9 +353,9 @@ __kernel void m17700_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m17800_a1-optimized.cl b/OpenCL/m17800_a1-optimized.cl index ce659a46c..2315e82a7 100644 --- a/OpenCL/m17800_a1-optimized.cl +++ b/OpenCL/m17800_a1-optimized.cl @@ -56,7 +56,7 @@ __kernel void m17800_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -64,9 +64,9 @@ __kernel void m17800_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -333,7 +333,7 @@ __kernel void m17800_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -353,9 +353,9 @@ __kernel void m17800_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m17900_a1-optimized.cl b/OpenCL/m17900_a1-optimized.cl index fec2665ca..978a6661e 100644 --- a/OpenCL/m17900_a1-optimized.cl +++ b/OpenCL/m17900_a1-optimized.cl @@ -56,7 +56,7 @@ __kernel void m17900_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -64,9 +64,9 @@ __kernel void m17900_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -333,7 +333,7 @@ __kernel void m17900_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -353,9 +353,9 @@ __kernel void m17900_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m18000_a1-optimized.cl b/OpenCL/m18000_a1-optimized.cl index 67e6af0c7..881850d43 100644 --- a/OpenCL/m18000_a1-optimized.cl +++ b/OpenCL/m18000_a1-optimized.cl @@ -56,7 +56,7 @@ __kernel void m18000_m04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * loop @@ -64,9 +64,9 @@ __kernel void m18000_m04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -333,7 +333,7 @@ __kernel void m18000_s04 (KERN_ATTR_BASIC ()) pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * digest @@ -353,9 +353,9 @@ __kernel void m18000_s04 (KERN_ATTR_BASIC ()) for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate diff --git a/OpenCL/m18200_a1-optimized.cl b/OpenCL/m18200_a1-optimized.cl index 69f570fbe..7de38dae1 100644 --- a/OpenCL/m18200_a1-optimized.cl +++ b/OpenCL/m18200_a1-optimized.cl @@ -589,7 +589,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -616,9 +616,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_m04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate @@ -736,7 +736,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s04 (KERN_A pw_buf1[2] = pws[gid].i[6]; pw_buf1[3] = pws[gid].i[7]; - const u32 pw_l_len = pws[gid].pw_len; + const u32 pw_l_len = pws[gid].pw_len & 63; /** * shared @@ -763,9 +763,9 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_s04 (KERN_A for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE) { - const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos); + const u32x pw_r_len = pwlenx_create_combt (combs_buf, il_pos) & 63; - const u32x pw_len = pw_l_len + pw_r_len; + const u32x pw_len = (pw_l_len + pw_r_len) & 63; /** * concat password candidate