arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-21 23:58:07 +00:00
Code Issues Releases Wiki Activity

Fixed maximum password length limit which was announced as 256 but actually was 255

Browse Source
This commit is contained in:
jsteube 2019-02-26 21:20:07 +01:00
parent f63762e597
commit 5da1e4b872
425 changed files with 775 additions and 776 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
OpenCL/inc_rp.h
View File

@ -62,12 +62,10 @@
#ifdef REAL_SHM
#define COPY_PW(x) \
__local pw_t s_pws[64]; \
s_pws[get_local_id(0)] = (x); \
s_pws[get_local_id(0)].pw_len &= 255;
s_pws[get_local_id(0)] = (x);
#else
#define COPY_PW(x) \
pw_t pw = (x); \
pw.pw_len &= 255;
pw_t pw = (x);
#endif
#ifdef REAL_SHM

8
OpenCL/m00000_a1-pure.cl
View File

@ -32,7 +32,7 @@ __kernel void m00000_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -42,7 +42,7 @@ __kernel void m00000_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx);
@ -86,7 +86,7 @@ __kernel void m00000_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -96,7 +96,7 @@ __kernel void m00000_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx);

4
OpenCL/m00000_a3-pure.cl
View File

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

8
OpenCL/m00010_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m00010_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m00010_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx, s, salt_len);
@ -106,7 +106,7 @@ __kernel void m00010_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -116,7 +116,7 @@ __kernel void m00010_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx, s, salt_len);

4
OpenCL/m00010_a3-pure.cl
View File

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

8
OpenCL/m00020_a1-pure.cl
View File

@ -34,7 +34,7 @@ __kernel void m00020_mxx (KERN_ATTR_BASIC ())
md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -44,7 +44,7 @@ __kernel void m00020_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx);
@ -90,7 +90,7 @@ __kernel void m00020_sxx (KERN_ATTR_BASIC ())
md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -100,7 +100,7 @@ __kernel void m00020_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx);

4
OpenCL/m00020_a3-pure.cl
View File

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

8
OpenCL/m00030_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m00030_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m00030_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx, s, salt_len);
@ -106,7 +106,7 @@ __kernel void m00030_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -116,7 +116,7 @@ __kernel void m00030_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx, s, salt_len);

4
OpenCL/m00030_a3-pure.cl
View File

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

8
OpenCL/m00040_a1-pure.cl
View File

@ -34,7 +34,7 @@ __kernel void m00040_mxx (KERN_ATTR_BASIC ())
md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -44,7 +44,7 @@ __kernel void m00040_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx);
@ -90,7 +90,7 @@ __kernel void m00040_sxx (KERN_ATTR_BASIC ())
md5_update_global (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -100,7 +100,7 @@ __kernel void m00040_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx);

4
OpenCL/m00040_a3-pure.cl
View File

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

4
OpenCL/m00050_a1-pure.cl
View File

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

4
OpenCL/m00050_a3-pure.cl
View File

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

4
OpenCL/m00060_a1-pure.cl
View File

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

4
OpenCL/m00060_a3-pure.cl
View File

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

8
OpenCL/m00100_a1-pure.cl
View File

@ -32,7 +32,7 @@ __kernel void m00100_mxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -42,7 +42,7 @@ __kernel void m00100_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx);
@ -86,7 +86,7 @@ __kernel void m00100_sxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -96,7 +96,7 @@ __kernel void m00100_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx);

4
OpenCL/m00100_a3-pure.cl
View File

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

8
OpenCL/m00110_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m00110_mxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m00110_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update (&ctx, s, salt_len);
@ -106,7 +106,7 @@ __kernel void m00110_sxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -116,7 +116,7 @@ __kernel void m00110_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update (&ctx, s, salt_len);

4
OpenCL/m00110_a3-pure.cl
View File

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

8
OpenCL/m00120_a1-pure.cl
View File

@ -34,7 +34,7 @@ __kernel void m00120_mxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -44,7 +44,7 @@ __kernel void m00120_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx);
@ -90,7 +90,7 @@ __kernel void m00120_sxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -100,7 +100,7 @@ __kernel void m00120_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx);

4
OpenCL/m00120_a3-pure.cl
View File

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

8
OpenCL/m00130_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m00130_mxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m00130_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update (&ctx, s, salt_len);
@ -106,7 +106,7 @@ __kernel void m00130_sxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -116,7 +116,7 @@ __kernel void m00130_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update (&ctx, s, salt_len);

4
OpenCL/m00130_a3-pure.cl
View File

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

8
OpenCL/m00140_a1-pure.cl
View File

@ -34,7 +34,7 @@ __kernel void m00140_mxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -44,7 +44,7 @@ __kernel void m00140_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx);
@ -90,7 +90,7 @@ __kernel void m00140_sxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -100,7 +100,7 @@ __kernel void m00140_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx);

4
OpenCL/m00140_a3-pure.cl
View File

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

4
OpenCL/m00150_a1-pure.cl
View File

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

4
OpenCL/m00150_a3-pure.cl
View File

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

4
OpenCL/m00160_a1-pure.cl
View File

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

4
OpenCL/m00160_a3-pure.cl
View File

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

8
OpenCL/m00300_a1-pure.cl
View File

@ -32,7 +32,7 @@ __kernel void m00300_mxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -42,7 +42,7 @@ __kernel void m00300_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx1 = ctx0;
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx1);
@ -109,7 +109,7 @@ __kernel void m00300_sxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -119,7 +119,7 @@ __kernel void m00300_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx1 = ctx0;
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx1);

4
OpenCL/m00300_a3-pure.cl
View File

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

4
OpenCL/m00400-pure.cl
View File

@ -36,7 +36,7 @@ __kernel void m00400_init (KERN_ATTR_TMPS (phpass_tmp_t))
md5_update_global (&md5_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len);
md5_final (&md5_ctx);
@ -67,7 +67,7 @@ __kernel void m00400_loop (KERN_ATTR_TMPS (phpass_tmp_t))
* init
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32 w[64] = { 0 };

4
OpenCL/m00500-pure.cl
View File

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

8
OpenCL/m00900_a1-pure.cl
View File

@ -32,7 +32,7 @@ __kernel void m00900_mxx (KERN_ATTR_BASIC ())
md4_init (&ctx0);
md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -42,7 +42,7 @@ __kernel void m00900_mxx (KERN_ATTR_BASIC ())
{
md4_ctx_t ctx = ctx0;
md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_final (&ctx);
@ -86,7 +86,7 @@ __kernel void m00900_sxx (KERN_ATTR_BASIC ())
md4_init (&ctx0);
md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -96,7 +96,7 @@ __kernel void m00900_sxx (KERN_ATTR_BASIC ())
{
md4_ctx_t ctx = ctx0;
md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_final (&ctx);

4
OpenCL/m00900_a3-pure.cl
View File

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

8
OpenCL/m01000_a1-pure.cl
View File

@ -32,7 +32,7 @@ __kernel void m01000_mxx (KERN_ATTR_BASIC ())
md4_init (&ctx0);
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -42,7 +42,7 @@ __kernel void m01000_mxx (KERN_ATTR_BASIC ())
{
md4_ctx_t ctx = ctx0;
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_final (&ctx);
@ -86,7 +86,7 @@ __kernel void m01000_sxx (KERN_ATTR_BASIC ())
md4_init (&ctx0);
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -96,7 +96,7 @@ __kernel void m01000_sxx (KERN_ATTR_BASIC ())
{
md4_ctx_t ctx = ctx0;
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_final (&ctx);

4
OpenCL/m01000_a3-pure.cl
View File

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

8
OpenCL/m01100_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m01100_mxx (KERN_ATTR_BASIC ())
md4_init (&ctx0);
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m01100_mxx (KERN_ATTR_BASIC ())
{
md4_ctx_t ctx1 = ctx0;
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_final (&ctx1);
@ -119,7 +119,7 @@ __kernel void m01100_sxx (KERN_ATTR_BASIC ())
md4_init (&ctx0);
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -129,7 +129,7 @@ __kernel void m01100_sxx (KERN_ATTR_BASIC ())
{
md4_ctx_t ctx1 = ctx0;
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_final (&ctx1);

4
OpenCL/m01100_a3-pure.cl
View File

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

8
OpenCL/m01300_a1-pure.cl
View File

@ -32,7 +32,7 @@ __kernel void m01300_mxx (KERN_ATTR_BASIC ())
sha224_init (&ctx0);
sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -42,7 +42,7 @@ __kernel void m01300_mxx (KERN_ATTR_BASIC ())
{
sha224_ctx_t ctx = ctx0;
sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha224_final (&ctx);
@ -86,7 +86,7 @@ __kernel void m01300_sxx (KERN_ATTR_BASIC ())
sha224_init (&ctx0);
sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -96,7 +96,7 @@ __kernel void m01300_sxx (KERN_ATTR_BASIC ())
{
sha224_ctx_t ctx = ctx0;
sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha224_final (&ctx);

4
OpenCL/m01300_a3-pure.cl
View File

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

8
OpenCL/m01400_a1-pure.cl
View File

@ -32,7 +32,7 @@ __kernel void m01400_mxx (KERN_ATTR_BASIC ())
sha256_init (&ctx0);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -42,7 +42,7 @@ __kernel void m01400_mxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_final (&ctx);
@ -86,7 +86,7 @@ __kernel void m01400_sxx (KERN_ATTR_BASIC ())
sha256_init (&ctx0);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -96,7 +96,7 @@ __kernel void m01400_sxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_final (&ctx);

4
OpenCL/m01400_a3-pure.cl
View File

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

8
OpenCL/m01410_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m01410_mxx (KERN_ATTR_BASIC ())
sha256_init (&ctx0);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m01410_mxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_update (&ctx, s, salt_len);
@ -106,7 +106,7 @@ __kernel void m01410_sxx (KERN_ATTR_BASIC ())
sha256_init (&ctx0);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -116,7 +116,7 @@ __kernel void m01410_sxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_update (&ctx, s, salt_len);

4
OpenCL/m01410_a3-pure.cl
View File

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

8
OpenCL/m01420_a1-pure.cl
View File

@ -34,7 +34,7 @@ __kernel void m01420_mxx (KERN_ATTR_BASIC ())
sha256_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -44,7 +44,7 @@ __kernel void m01420_mxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_final (&ctx);
@ -90,7 +90,7 @@ __kernel void m01420_sxx (KERN_ATTR_BASIC ())
sha256_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -100,7 +100,7 @@ __kernel void m01420_sxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_final (&ctx);

4
OpenCL/m01420_a3-pure.cl
View File

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

8
OpenCL/m01430_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m01430_mxx (KERN_ATTR_BASIC ())
sha256_init (&ctx0);
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m01430_mxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_update (&ctx, s, salt_len);
@ -106,7 +106,7 @@ __kernel void m01430_sxx (KERN_ATTR_BASIC ())
sha256_init (&ctx0);
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -116,7 +116,7 @@ __kernel void m01430_sxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_update (&ctx, s, salt_len);

4
OpenCL/m01430_a3-pure.cl
View File

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

8
OpenCL/m01440_a1-pure.cl
View File

@ -34,7 +34,7 @@ __kernel void m01440_mxx (KERN_ATTR_BASIC ())
sha256_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -44,7 +44,7 @@ __kernel void m01440_mxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_final (&ctx);
@ -90,7 +90,7 @@ __kernel void m01440_sxx (KERN_ATTR_BASIC ())
sha256_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -100,7 +100,7 @@ __kernel void m01440_sxx (KERN_ATTR_BASIC ())
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_final (&ctx);

4
OpenCL/m01440_a3-pure.cl
View File

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

4
OpenCL/m01450_a1-pure.cl
View File

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

4
OpenCL/m01450_a3-pure.cl
View File

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

4
OpenCL/m01460_a1-pure.cl
View File

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

4
OpenCL/m01460_a3-pure.cl
View File

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

4
OpenCL/m01600-pure.cl
View File

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

8
OpenCL/m01700_a1-pure.cl
View File

@ -32,7 +32,7 @@ __kernel void m01700_mxx (KERN_ATTR_BASIC ())
sha512_init (&ctx0);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -42,7 +42,7 @@ __kernel void m01700_mxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_final (&ctx);
@ -86,7 +86,7 @@ __kernel void m01700_sxx (KERN_ATTR_BASIC ())
sha512_init (&ctx0);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -96,7 +96,7 @@ __kernel void m01700_sxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_final (&ctx);

4
OpenCL/m01700_a3-pure.cl
View File

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

8
OpenCL/m01710_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m01710_mxx (KERN_ATTR_BASIC ())
sha512_init (&ctx0);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m01710_mxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_update (&ctx, s, salt_len);
@ -106,7 +106,7 @@ __kernel void m01710_sxx (KERN_ATTR_BASIC ())
sha512_init (&ctx0);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -116,7 +116,7 @@ __kernel void m01710_sxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_update (&ctx, s, salt_len);

4
OpenCL/m01710_a3-pure.cl
View File

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

8
OpenCL/m01720_a1-pure.cl
View File

@ -34,7 +34,7 @@ __kernel void m01720_mxx (KERN_ATTR_BASIC ())
sha512_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -44,7 +44,7 @@ __kernel void m01720_mxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_final (&ctx);
@ -90,7 +90,7 @@ __kernel void m01720_sxx (KERN_ATTR_BASIC ())
sha512_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -100,7 +100,7 @@ __kernel void m01720_sxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_final (&ctx);

4
OpenCL/m01720_a3-pure.cl
View File

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

8
OpenCL/m01730_a1-pure.cl
View File

@ -41,7 +41,7 @@ __kernel void m01730_mxx (KERN_ATTR_BASIC ())
sha512_init (&ctx0);
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -51,7 +51,7 @@ __kernel void m01730_mxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_update (&ctx, s, salt_len);
@ -106,7 +106,7 @@ __kernel void m01730_sxx (KERN_ATTR_BASIC ())
sha512_init (&ctx0);
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -116,7 +116,7 @@ __kernel void m01730_sxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_update (&ctx, s, salt_len);

4
OpenCL/m01730_a3-pure.cl
View File

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

8
OpenCL/m01740_a1-pure.cl
View File

@ -34,7 +34,7 @@ __kernel void m01740_mxx (KERN_ATTR_BASIC ())
sha512_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -44,7 +44,7 @@ __kernel void m01740_mxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_final (&ctx);
@ -90,7 +90,7 @@ __kernel void m01740_sxx (KERN_ATTR_BASIC ())
sha512_update_global_swap (&ctx0, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -100,7 +100,7 @@ __kernel void m01740_sxx (KERN_ATTR_BASIC ())
{
sha512_ctx_t ctx = ctx0;
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha512_final (&ctx);

4
OpenCL/m01740_a3-pure.cl
View File

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

4
OpenCL/m01750_a1-pure.cl
View File

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

4
OpenCL/m01750_a3-pure.cl
View File

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

4
OpenCL/m01760_a1-pure.cl
View File

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

4
OpenCL/m01760_a3-pure.cl
View File

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

4
OpenCL/m01800-pure.cl
View File

@ -27,7 +27,7 @@ __kernel void m01800_init (KERN_ATTR_TMPS (sha512crypt_tmp_t))
* init
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32 w[64] = { 0 };
@ -308,7 +308,7 @@ __kernel void m01800_loop (KERN_ATTR_TMPS (sha512crypt_tmp_t))
if (gid >= gid_max) return;
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
const u32 salt_len = salt_bufs[salt_pos].salt_len;

2
OpenCL/m02100-pure.cl
View File

@ -69,7 +69,7 @@ __kernel void m02100_init (KERN_ATTR_TMPS (dcc2_tmp_t))
md4_init (&md4_ctx1);
md4_update_global_utf16le (&md4_ctx1, pws[gid].i, pws[gid].pw_len & 255);
md4_update_global_utf16le (&md4_ctx1, pws[gid].i, pws[gid].pw_len);
md4_final (&md4_ctx1);

2
OpenCL/m02500-pure.cl
View File

@ -97,7 +97,7 @@ __kernel void m02500_init (KERN_ATTR_TMPS_ESALT (wpa_pbkdf2_tmp_t, wpa_eapol_t))
sha1_hmac_ctx_t sha1_hmac_ctx;
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];

8
OpenCL/m02610_a1-pure.cl
View File

@ -71,7 +71,7 @@ __kernel void m02610_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -81,7 +81,7 @@ __kernel void m02610_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);
@ -184,7 +184,7 @@ __kernel void m02610_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -194,7 +194,7 @@ __kernel void m02610_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);

4
OpenCL/m02610_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m02610_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -184,7 +184,7 @@ __kernel void m02610_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m02810_a1-pure.cl
View File

@ -71,7 +71,7 @@ __kernel void m02810_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -81,7 +81,7 @@ __kernel void m02810_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);
@ -207,7 +207,7 @@ __kernel void m02810_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -217,7 +217,7 @@ __kernel void m02810_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);

4
OpenCL/m02810_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m02810_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -207,7 +207,7 @@ __kernel void m02810_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

2
OpenCL/m03200-pure.cl
View File

@ -359,7 +359,7 @@ __kernel void __attribute__((reqd_work_group_size(8, 1, 1))) m03200_init (KERN_A
if (gid >= gid_max) return;
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32 w[18];

8
OpenCL/m03710_a1-pure.cl
View File

@ -71,7 +71,7 @@ __kernel void m03710_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -81,7 +81,7 @@ __kernel void m03710_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);
@ -197,7 +197,7 @@ __kernel void m03710_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -207,7 +207,7 @@ __kernel void m03710_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);

4
OpenCL/m03710_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m03710_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -197,7 +197,7 @@ __kernel void m03710_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m03800_a1-pure.cl
View File

@ -43,7 +43,7 @@ __kernel void m03800_mxx (KERN_ATTR_BASIC ())
md5_update (&ctx0, s, salt_len);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -53,7 +53,7 @@ __kernel void m03800_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx, s, salt_len);
@ -110,7 +110,7 @@ __kernel void m03800_sxx (KERN_ATTR_BASIC ())
md5_update (&ctx0, s, salt_len);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -120,7 +120,7 @@ __kernel void m03800_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx, s, salt_len);

4
OpenCL/m03800_a3-pure.cl
View File

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

8
OpenCL/m03910_a1-pure.cl
View File

@ -71,7 +71,7 @@ __kernel void m03910_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -81,7 +81,7 @@ __kernel void m03910_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);
@ -207,7 +207,7 @@ __kernel void m03910_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -217,7 +217,7 @@ __kernel void m03910_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);

4
OpenCL/m03910_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m03910_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -207,7 +207,7 @@ __kernel void m03910_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m04010_a1-pure.cl
View File

@ -66,7 +66,7 @@ __kernel void m04010_mxx (KERN_ATTR_BASIC ())
md5_ctx_t ctx0t = ctx0;
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -76,7 +76,7 @@ __kernel void m04010_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0t;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);
@ -183,7 +183,7 @@ __kernel void m04010_sxx (KERN_ATTR_BASIC ())
md5_ctx_t ctx0t = ctx0;
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -193,7 +193,7 @@ __kernel void m04010_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0t;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);

4
OpenCL/m04010_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m04010_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -192,7 +192,7 @@ __kernel void m04010_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m04110_a1-pure.cl
View File

@ -77,7 +77,7 @@ __kernel void m04110_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0t);
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -87,7 +87,7 @@ __kernel void m04110_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0t;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx1, s, salt_len);
@ -207,7 +207,7 @@ __kernel void m04110_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0t);
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -217,7 +217,7 @@ __kernel void m04110_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0t;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx1, s, salt_len);

4
OpenCL/m04110_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m04110_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -203,7 +203,7 @@ __kernel void m04110_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m04310_a1-pure.cl
View File

@ -71,7 +71,7 @@ __kernel void m04310_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -81,7 +81,7 @@ __kernel void m04310_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);
@ -184,7 +184,7 @@ __kernel void m04310_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -194,7 +194,7 @@ __kernel void m04310_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);

4
OpenCL/m04310_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m04310_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -184,7 +184,7 @@ __kernel void m04310_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m04400_a1-pure.cl
View File

@ -63,7 +63,7 @@ __kernel void m04400_mxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -73,7 +73,7 @@ __kernel void m04400_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx1 = ctx0;
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx1);
@ -170,7 +170,7 @@ __kernel void m04400_sxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -180,7 +180,7 @@ __kernel void m04400_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx1 = ctx0;
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx1);

4
OpenCL/m04400_a3-pure.cl
View File

@ -59,7 +59,7 @@ __kernel void m04400_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -179,7 +179,7 @@ __kernel void m04400_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m04500_a1-pure.cl
View File

@ -62,7 +62,7 @@ __kernel void m04500_mxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -72,7 +72,7 @@ __kernel void m04500_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx1 = ctx0;
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx1);
@ -169,7 +169,7 @@ __kernel void m04500_sxx (KERN_ATTR_BASIC ())
sha1_init (&ctx0);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -179,7 +179,7 @@ __kernel void m04500_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx1 = ctx0;
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx1);

4
OpenCL/m04500_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m04500_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -178,7 +178,7 @@ __kernel void m04500_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m04520_a1-pure.cl
View File

@ -68,7 +68,7 @@ __kernel void m04520_mxx (KERN_ATTR_BASIC ())
sha1_init (&ctx1l);
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -78,7 +78,7 @@ __kernel void m04520_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx1 = ctx1l;
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx1);
@ -190,7 +190,7 @@ __kernel void m04520_sxx (KERN_ATTR_BASIC ())
sha1_init (&ctx1l);
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -200,7 +200,7 @@ __kernel void m04520_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx1 = ctx1l;
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_final (&ctx1);

4
OpenCL/m04520_a3-pure.cl
View File

@ -58,7 +58,7 @@ __kernel void m04520_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -195,7 +195,7 @@ __kernel void m04520_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m04700_a1-pure.cl
View File

@ -63,7 +63,7 @@ __kernel void m04700_mxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -73,7 +73,7 @@ __kernel void m04700_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);
@ -165,7 +165,7 @@ __kernel void m04700_sxx (KERN_ATTR_BASIC ())
md5_init (&ctx0);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -175,7 +175,7 @@ __kernel void m04700_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_final (&ctx1);

4
OpenCL/m04700_a3-pure.cl
View File

@ -59,7 +59,7 @@ __kernel void m04700_mxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };
@ -174,7 +174,7 @@ __kernel void m04700_sxx (KERN_ATTR_VECTOR ())
* base
*/
const u32 pw_len = pws[gid].pw_len & 255;
const u32 pw_len = pws[gid].pw_len;
u32x w[64] = { 0 };

8
OpenCL/m04800_a1-pure.cl
View File

@ -45,7 +45,7 @@ __kernel void m04800_mxx (KERN_ATTR_BASIC ())
ctx0.len = 1;
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -55,7 +55,7 @@ __kernel void m04800_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx, s, salt_len);
@ -114,7 +114,7 @@ __kernel void m04800_sxx (KERN_ATTR_BASIC ())
ctx0.len = 1;
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
/**
* loop
@ -124,7 +124,7 @@ __kernel void m04800_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update (&ctx, s, salt_len);

4
OpenCL/m04800_a3-pure.cl
View File

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

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