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

Get rid of ceil() in OpenCL kernels

Browse Source
pull/1304/head
jsteube 7 years ago
parent 53f53fe014
commit ad1ce462d1
154 changed files with 401 additions and 1174 deletions
Show Stats Download Patch File Download Diff File

48
OpenCL/inc_rp.cl
Unescape View File

@ -163,13 +163,11 @@ static void exchange_byte (u32 *buf, const int off_src, const int off_dst)
static int mangle_lrest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
{
const int lenv = ceil ((float) len / 4);
for (int i = 0; i < lenv; i++)
for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
{
const u32 t = buf[i];
const u32 t = buf[idx];
buf[i] = t | generate_cmask (t);
buf[idx] = t | generate_cmask (t);
}
return (len);
@ -177,13 +175,11 @@ static int mangle_lrest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32
static int mangle_lrest_ufirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
{
const int lenv = ceil ((float) len / 4);
for (int i = 0; i < lenv; i++)
for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
{
const u32 t = buf[i];
const u32 t = buf[idx];
buf[i] = t | generate_cmask (t);
buf[idx] = t | generate_cmask (t);
}
const u32 t = buf[0];
@ -195,13 +191,11 @@ static int mangle_lrest_ufirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8
static int mangle_urest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
{
const int lenv = ceil ((float) len / 4);
for (int i = 0; i < lenv; i++)
for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
{
const u32 t = buf[i];
const u32 t = buf[idx];
buf[i] = t & ~(generate_cmask (t));
buf[idx] = t & ~(generate_cmask (t));
}
return (len);
@ -209,13 +203,11 @@ static int mangle_urest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32
static int mangle_urest_lfirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
{
const int lenv = ceil ((float) len / 4);
for (int i = 0; i < lenv; i++)
for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
{
const u32 t = buf[i];
const u32 t = buf[idx];
buf[i] = t & ~(generate_cmask (t));
buf[idx] = t & ~(generate_cmask (t));
}
const u32 t = buf[0];
@ -227,13 +219,11 @@ static int mangle_urest_lfirst (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8
static int mangle_trest (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1, u32 *buf, const int len)
{
const int lenv = ceil ((float) len / 4);
for (int i = 0; i < lenv; i++)
for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
{
const u32 t = buf[i];
const u32 t = buf[idx];
buf[i] = t ^ generate_cmask (t);
buf[idx] = t ^ generate_cmask (t);
}
return (len);
@ -668,9 +658,9 @@ static int mangle_title_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1,
mangle_lrest_ufirst (0, 0, buf, len);
for (int i = 0; i < ceil ((float) len / 4); i++)
for (int i = 0, idx = 0; i < len; i += 4, idx += 1)
{
const u32 v = buf[i];
const u32 v = buf[idx];
u32 out0 = 0;
u32 out1 = 0;
@ -680,8 +670,8 @@ static int mangle_title_sep (MAYBE_UNUSED const u8 p0, MAYBE_UNUSED const u8 p1,
if (((v >> 16) & 0xff) == p0) out0 |= 0xff000000;
if (((v >> 24) & 0xff) == p0) out1 |= 0x000000ff;
buf[i + 0] &= ~(generate_cmask (buf[i + 0]) & out0);
buf[i + 1] &= ~(generate_cmask (buf[i + 1]) & out1);
buf[idx + 0] &= ~(generate_cmask (buf[idx + 0]) & out0);
buf[idx + 1] &= ~(generate_cmask (buf[idx + 1]) & out1);
}
return (len);

8
OpenCL/m00000_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -99,11 +97,9 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m00010_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -103,11 +101,9 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m00010_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -97,11 +95,9 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m00010_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -112,22 +108,18 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m00020_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -105,11 +103,9 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m00030_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -103,11 +101,9 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m00030_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -97,11 +95,9 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m00030_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -112,22 +108,18 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m00040_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -105,11 +103,9 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m00050_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -101,11 +99,9 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m00050_a1.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -124,22 +120,18 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m00050_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -110,22 +106,18 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m00060_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -103,11 +101,9 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m00060_a1.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -126,22 +122,18 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m00060_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -112,22 +108,18 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m00100_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -99,11 +97,9 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m00110_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m00110_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -97,11 +95,9 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m00110_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -112,22 +108,18 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m00120_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -105,11 +103,9 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m00130_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m00130_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -97,11 +95,9 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

17
OpenCL/m00130_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -83,7 +79,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
}
}
__kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -113,22 +108,18 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m00140_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -105,11 +103,9 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m00150_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -101,11 +99,9 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m00150_a1.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -124,22 +120,18 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m00150_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -110,22 +106,18 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m00160_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m00160_a1.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -126,22 +122,18 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m00160_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -112,22 +108,18 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m00300_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -122,11 +120,9 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

4
OpenCL/m00400.cl
Unescape View File

@ -69,11 +69,9 @@ __kernel void m00400_loop (__global pw_t *pws, __global const kernel_rule_t *rul
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

16
OpenCL/m00500.cl
Unescape View File

@ -33,22 +33,18 @@ __kernel void m00500_init (__global pw_t *pws, __global const kernel_rule_t *rul
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -143,22 +139,18 @@ __kernel void m00500_loop (__global pw_t *pws, __global const kernel_rule_t *rul
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m00900_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -99,11 +97,9 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01000_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -99,11 +97,9 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01100_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -116,11 +114,9 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m01100_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -110,11 +108,9 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m01100_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -125,22 +121,18 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m01300_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -99,11 +97,9 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01400_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -99,11 +97,9 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01410_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01410_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -97,11 +95,9 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01410_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -112,22 +108,18 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01420_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -105,11 +103,9 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01430_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01430_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -97,11 +95,9 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

17
OpenCL/m01430_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -83,7 +79,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
}
}
__kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -113,22 +108,18 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01440_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -105,11 +103,9 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01450_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -101,11 +99,9 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01450_a1.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -124,22 +120,18 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01450_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -110,22 +106,18 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01460_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01460_a1.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -126,22 +122,18 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01460_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -112,22 +108,18 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01600.cl
Unescape View File

@ -34,22 +34,18 @@ __kernel void m01600_init (__global pw_t *pws, __global const kernel_rule_t *rul
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -145,22 +141,18 @@ __kernel void m01600_loop (__global pw_t *pws, __global const kernel_rule_t *rul
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m01700_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -99,11 +97,9 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01710_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01710_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -97,11 +95,9 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01710_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -112,22 +108,18 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01720_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -105,11 +103,9 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01730_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01730_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -97,11 +95,9 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

17
OpenCL/m01730_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -83,7 +79,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
}
}
__kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __global void *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max)
{
/**
@ -113,22 +108,18 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01740_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -105,11 +103,9 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m01750_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -101,11 +99,9 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01750_a1.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -124,22 +120,18 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01750_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -110,22 +106,18 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

8
OpenCL/m01760_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -103,11 +101,9 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01760_a1.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -126,22 +122,18 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32_S (pws[gid].i[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

16
OpenCL/m01760_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m01760_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -112,22 +108,18 @@ __kernel void m01760_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

12
OpenCL/m01800.cl
Unescape View File

@ -29,32 +29,28 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32 w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = swap32 (w[idx]);
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32 (s[idx]);
}

8
OpenCL/m02610_a0.cl
Unescape View File

@ -64,11 +64,9 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -181,11 +179,9 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m02610_a1.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -175,11 +173,9 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m02610_a3.cl
Unescape View File

@ -60,22 +60,18 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -190,22 +186,18 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m02810_a0.cl
Unescape View File

@ -64,11 +64,9 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}
@ -204,11 +202,9 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}

8
OpenCL/m02810_a1.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}
@ -198,11 +196,9 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}

16
OpenCL/m02810_a3.cl
Unescape View File

@ -60,22 +60,18 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}
@ -213,22 +209,18 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}

8
OpenCL/m03710_a0.cl
Unescape View File

@ -64,11 +64,9 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -194,11 +192,9 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m03710_a1.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -188,11 +186,9 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m03710_a3.cl
Unescape View File

@ -60,22 +60,18 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -203,22 +199,18 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m03800_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -107,11 +105,9 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m03800_a1.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -99,11 +97,9 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m03800_a3.cl
Unescape View File

@ -30,22 +30,18 @@ __kernel void m03800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -118,22 +114,18 @@ __kernel void m03800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m03910_a0.cl
Unescape View File

@ -64,11 +64,9 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}
@ -204,11 +202,9 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}

8
OpenCL/m03910_a1.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}
@ -198,11 +196,9 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}

16
OpenCL/m03910_a3.cl
Unescape View File

@ -60,22 +60,18 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}
@ -213,22 +209,18 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = 32;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[8] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf_pc[idx];
}

8
OpenCL/m04010_a3.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m04010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -196,11 +194,9 @@ __kernel void m04010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m04110_a0.cl
Unescape View File

@ -64,11 +64,9 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -198,11 +196,9 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m04110_a1.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -192,11 +190,9 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m04110_a3.cl
Unescape View File

@ -60,22 +60,18 @@ __kernel void m04110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -209,22 +205,18 @@ __kernel void m04110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m04310_a0.cl
Unescape View File

@ -64,11 +64,9 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -181,11 +179,9 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m04310_a1.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -175,11 +173,9 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

16
OpenCL/m04310_a3.cl
Unescape View File

@ -60,22 +60,18 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}
@ -190,22 +186,18 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32x s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
}

8
OpenCL/m04400_a3.cl
Unescape View File

@ -61,11 +61,9 @@ __kernel void m04400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -183,11 +181,9 @@ __kernel void m04400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m04500_a3.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m04500_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -182,11 +180,9 @@ __kernel void m04500_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m04520_a3.cl
Unescape View File

@ -60,11 +60,9 @@ __kernel void m04520_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -199,11 +197,9 @@ __kernel void m04520_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m04700_a3.cl
Unescape View File

@ -61,11 +61,9 @@ __kernel void m04700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -178,11 +176,9 @@ __kernel void m04700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m04800_a3.cl
Unescape View File

@ -30,11 +30,9 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}
@ -118,11 +116,9 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 pw_len = pws[gid].pw_len;
const u32 pw_lenv = ceil ((float) pw_len / 4);
u32x w[64] = { 0 };
for (int idx = 0; idx < pw_lenv; idx++)
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
{
w[idx] = pws[gid].i[idx];
}

8
OpenCL/m04900_a0.cl
Unescape View File

@ -34,11 +34,9 @@ __kernel void m04900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}
@ -107,11 +105,9 @@ __kernel void m04900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
const u32 salt_len = salt_bufs[salt_pos].salt_len;
const u32 salt_lenv = ceil ((float) salt_len / 4);
u32 s[64] = { 0 };
for (int idx = 0; idx < salt_lenv; idx++)
for (int i = 0, idx = 0; i < salt_len; i += 4, idx += 1)
{
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
}

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

Write Preview
Loading…
Cancel
Save