arno
/
hashcat
mirror of https://github.com/hashcat/hashcat.git
1
0
Fork 0
Code Issues Releases Wiki Activity

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

Browse Source
pull/1793/head
Jens Steube 6 years ago
parent 53c8600089
commit 64dfd40113
121 changed files with 722 additions and 722 deletions
Show Stats Download Patch File Download Diff File

12
OpenCL/m00000_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00010_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00020_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00030_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00040_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00050_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00060_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00100_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00110_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00120_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00130_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00140_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00150_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00160_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00200_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m00300_a1-optimized.cl
Unescape View File

@ -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

8
OpenCL/m00600_a1-optimized.cl
Unescape View File

@ -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;

12
OpenCL/m00900_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01000_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01100_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01300_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01400_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01410_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01420_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01430_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01440_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01450_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01460_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01500_a1-pure.cl
Unescape View File

@ -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

12
OpenCL/m01700_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01710_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01720_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01730_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01740_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01750_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m01760_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m02400_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m02410_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m02610_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m02710_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m02810_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m03000_a1-pure.cl
Unescape View File

@ -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

12
OpenCL/m03100_a1-optimized.cl
Unescape View File

@ -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;

12
OpenCL/m03710_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m03800_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m03910_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m04010_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m04110_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m04310_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m04400_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m04500_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m04520_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m04700_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m04800_a1-optimized.cl
Unescape View File

@ -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;

12
OpenCL/m04900_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m05100_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m05300_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m05400_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m05500_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m05600_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m06000_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m06100_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m06900_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m07000_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m07300_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m07500_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m07700_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m07701_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m07800_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m07801_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m08000_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m08100_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m08300_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m08400_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m08500_a1-pure.cl
Unescape View File

@ -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

12
OpenCL/m08600_a1-pure.cl
Unescape View File

@ -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

12
OpenCL/m08700_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m09700_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m09710_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m09720_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m09800_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m09810_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m09820_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m09900_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m10100_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m10400_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m10410_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m10420_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m10800_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m11000_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m11100_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m11200_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m11500_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m11700_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m11800_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m12600_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m13100_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m13300_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m13500_a1-optimized.cl
Unescape View File

@ -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

12
OpenCL/m13800_a1-optimized.cl
Unescape View File

@ -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

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

Write Preview
Loading…
Cancel
Save