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

pull/1793/head
Jens Steube 6 years ago
parent 0e428b3c40
commit 53c8600089

@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);
@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);

@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx, s, salt_len);
@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx, s, salt_len);

@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);
@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);

@ -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);
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx, s, salt_len);
@ -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);
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx, s, salt_len);

@ -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);
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);
@ -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);
md5_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);

@ -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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);
@ -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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);

@ -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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx, s, salt_len);
@ -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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx, s, salt_len);

@ -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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);
@ -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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);

@ -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);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx, s, salt_len);
@ -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);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx, s, salt_len);

@ -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);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);
@ -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);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);

@ -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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);
@ -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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);

@ -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);
md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);
@ -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);
md4_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);

@ -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);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);
@ -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);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);

@ -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);
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx1);
@ -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);
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx1);

@ -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);
sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha224_final (&ctx);
@ -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);
sha224_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha224_final (&ctx);

@ -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);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_final (&ctx);
@ -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);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_final (&ctx);

@ -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);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update (&ctx, s, salt_len);
@ -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);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update (&ctx, s, salt_len);

@ -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);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_final (&ctx);
@ -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);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_final (&ctx);

@ -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);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update (&ctx, s, salt_len);
@ -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);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update (&ctx, s, salt_len);

@ -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);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_final (&ctx);
@ -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);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_final (&ctx);

@ -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);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_final (&ctx);
@ -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);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_final (&ctx);

@ -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);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update (&ctx, s, salt_len);
@ -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);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update (&ctx, s, salt_len);

@ -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);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_final (&ctx);
@ -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);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_final (&ctx);

@ -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);
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update (&ctx, s, salt_len);
@ -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);
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update (&ctx, s, salt_len);

@ -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);
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_final (&ctx);
@ -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);
sha512_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_final (&ctx);

@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);
@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);

@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);
@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);

@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);
@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);

@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx, s, salt_len);
@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx, s, salt_len);

@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);
@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);

@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);
@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);

@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx1, s, salt_len);
@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx1, s, salt_len);

@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);
@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);

@ -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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);
@ -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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);

@ -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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);
@ -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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);

@ -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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);
@ -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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);

@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);
@ -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);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);

@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx, s, salt_len);
@ -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);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_update (&ctx, s, salt_len);

@ -53,7 +53,7 @@ __kernel void m04900_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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx, s, salt_len);
@ -120,7 +120,7 @@ __kernel void m04900_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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx, s, salt_len);

@ -42,7 +42,7 @@ __kernel void m05100_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);
@ -100,7 +100,7 @@ __kernel void m05100_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);

@ -555,7 +555,7 @@ __kernel void m05500_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);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);
@ -689,7 +689,7 @@ __kernel void m05500_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);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);

@ -43,7 +43,7 @@ __kernel void m05600_mxx (KERN_ATTR_ESALT (netntlm_t))
{
md4_ctx_t ctx1 = ctx10;
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx1);
@ -152,7 +152,7 @@ __kernel void m05600_sxx (KERN_ATTR_ESALT (netntlm_t))
{
md4_ctx_t ctx1 = ctx10;
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_update_global_utf16le (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx1);

@ -42,7 +42,7 @@ __kernel void m06000_mxx (KERN_ATTR_BASIC ())
{
ripemd160_ctx_t ctx = ctx0;
ripemd160_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
ripemd160_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
ripemd160_final (&ctx);
@ -96,7 +96,7 @@ __kernel void m06000_sxx (KERN_ATTR_BASIC ())
{
ripemd160_ctx_t ctx = ctx0;
ripemd160_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
ripemd160_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
ripemd160_final (&ctx);

@ -73,7 +73,7 @@ __kernel void m06100_mxx (KERN_ATTR_BASIC ())
{
whirlpool_ctx_t ctx = ctx0;
whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
whirlpool_final (&ctx);
@ -158,7 +158,7 @@ __kernel void m06100_sxx (KERN_ATTR_BASIC ())
{
whirlpool_ctx_t ctx = ctx0;
whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
whirlpool_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
whirlpool_final (&ctx);

@ -44,7 +44,7 @@ __kernel void m07000_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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
/**
* pepper
@ -128,7 +128,7 @@ __kernel void m07000_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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
/**
* pepper

@ -317,7 +317,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_A
{
md4_ctx_t ctx = ctx0;
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);
@ -386,7 +386,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_sxx (KERN_A
{
md4_ctx_t ctx = ctx0;
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);

@ -64,7 +64,7 @@ __kernel void m08300_mxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx1, s_pc, salt_len_pc + 1);
@ -179,7 +179,7 @@ __kernel void m08300_sxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx1, s_pc, salt_len_pc + 1);

@ -78,7 +78,7 @@ __kernel void m08400_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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);
@ -239,7 +239,7 @@ __kernel void m08400_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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);

@ -42,7 +42,7 @@ __kernel void m10800_mxx (KERN_ATTR_BASIC ())
{
sha384_ctx_t ctx = ctx0;
sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha384_final (&ctx);
@ -96,7 +96,7 @@ __kernel void m10800_sxx (KERN_ATTR_BASIC ())
{
sha384_ctx_t ctx = ctx0;
sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha384_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha384_final (&ctx);

@ -44,7 +44,7 @@ __kernel void m11000_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);
@ -100,7 +100,7 @@ __kernel void m11000_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx);

@ -94,7 +94,7 @@ __kernel void m11100_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0t;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
u32 s0[4];
u32 s1[4];
@ -250,7 +250,7 @@ __kernel void m11100_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx1 = ctx0t;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
u32 s0[4];
u32 s1[4];

@ -48,7 +48,7 @@ __kernel void m11200_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx2 = ctx2l;
sha1_update_global_swap (&ctx2, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_swap (&ctx2, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx2);
@ -174,7 +174,7 @@ __kernel void m11200_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx2 = ctx2l;
sha1_update_global_swap (&ctx2, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_swap (&ctx2, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx2);

@ -74,7 +74,7 @@ __kernel void m11400_mxx (KERN_ATTR_ESALT (sip_t))
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);
@ -180,7 +180,7 @@ __kernel void m11400_sxx (KERN_ATTR_ESALT (sip_t))
{
md5_ctx_t ctx1 = ctx0;
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md5_final (&ctx1);

@ -71,7 +71,7 @@ __kernel void m11700_mxx (KERN_ATTR_BASIC ())
{
streebog256_ctx_t ctx = ctx0;
streebog256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
streebog256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
streebog256_final (&ctx);
@ -154,7 +154,7 @@ __kernel void m11700_sxx (KERN_ATTR_BASIC ())
{
streebog256_ctx_t ctx = ctx0;
streebog256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
streebog256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
streebog256_final (&ctx);

@ -71,7 +71,7 @@ __kernel void m11800_mxx (KERN_ATTR_BASIC ())
{
streebog512_ctx_t ctx = ctx0;
streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
streebog512_final (&ctx);
@ -154,7 +154,7 @@ __kernel void m11800_sxx (KERN_ATTR_BASIC ())
{
streebog512_ctx_t ctx = ctx0;
streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
streebog512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
streebog512_final (&ctx);

@ -88,7 +88,7 @@ __kernel void m12600_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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);
@ -234,7 +234,7 @@ __kernel void m12600_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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);

@ -415,7 +415,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_A
{
md4_ctx_t ctx = ctx0;
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);
@ -475,7 +475,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_sxx (KERN_A
{
md4_ctx_t ctx = ctx0;
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);

@ -42,7 +42,7 @@ __kernel void m13300_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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);
@ -98,7 +98,7 @@ __kernel void m13300_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);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);

@ -71,7 +71,7 @@ __kernel void m13500_mxx (KERN_ATTR_ESALT (pstoken_t))
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);
@ -154,7 +154,7 @@ __kernel void m13500_sxx (KERN_ATTR_ESALT (pstoken_t))
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx);

@ -42,7 +42,7 @@ __kernel void m13800_mxx (KERN_ATTR_ESALT (win8phone_t))
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global (&ctx, esalt_bufs[digests_offset].salt_buf, 128);
@ -98,7 +98,7 @@ __kernel void m13800_sxx (KERN_ATTR_ESALT (win8phone_t))
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_update_global_utf16le_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_update_global (&ctx, esalt_bufs[digests_offset].salt_buf, 128);

@ -78,7 +78,7 @@ __kernel void m13900_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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);
@ -239,7 +239,7 @@ __kernel void m13900_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);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_final (&ctx1);

@ -141,7 +141,7 @@ __kernel void m14400_mxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
d40[0] = 0x2d2d2d2d;
d40[1] = 0;
@ -231,7 +231,7 @@ __kernel void m14400_mxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
d40[0] = 0x2d2d2d2d;
d40[1] = 0;
@ -398,7 +398,7 @@ __kernel void m14400_sxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_swap (&ctx1, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
d40[0] = 0x2d2d2d2d;
d40[1] = 0;
@ -488,7 +488,7 @@ __kernel void m14400_sxx (KERN_ATTR_BASIC ())
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
d40[0] = 0x2d2d2d2d;
d40[1] = 0;

@ -51,7 +51,7 @@ __kernel void m15000_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);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update (&ctx, s, salt_len);
@ -116,7 +116,7 @@ __kernel void m15000_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);
sha512_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha512_update (&ctx, s, salt_len);

@ -51,7 +51,7 @@ __kernel void m15500_mxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_utf16be_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_utf16be_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx, s, salt_len);
@ -122,7 +122,7 @@ __kernel void m15500_sxx (KERN_ATTR_BASIC ())
{
sha1_ctx_t ctx = ctx0;
sha1_update_global_utf16be_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha1_update_global_utf16be_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha1_update (&ctx, s, salt_len);

@ -74,7 +74,7 @@ __kernel void m16100_mxx (KERN_ATTR_ESALT (tacacs_plus_t))
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
u32 sequence0[4];
u32 sequence1[4];
@ -232,7 +232,7 @@ __kernel void m16100_sxx (KERN_ATTR_ESALT (tacacs_plus_t))
{
md5_ctx_t ctx = ctx0;
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
u32 sequence0[4];
u32 sequence1[4];

@ -204,7 +204,7 @@ __kernel void m16400_mxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
cram_md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
cram_md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
cram_md5_final (&ctx);
@ -258,7 +258,7 @@ __kernel void m16400_sxx (KERN_ATTR_BASIC ())
{
md5_ctx_t ctx = ctx0;
cram_md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
cram_md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
cram_md5_final (&ctx);

@ -95,7 +95,7 @@ __kernel void m16600_mxx (KERN_ATTR_ESALT (electrum_wallet_t))
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_final (&ctx);
@ -268,7 +268,7 @@ __kernel void m16600_sxx (KERN_ATTR_ESALT (electrum_wallet_t))
{
sha256_ctx_t ctx = ctx0;
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
sha256_update_global_swap (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
sha256_final (&ctx);

@ -413,7 +413,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_mxx (KERN_A
{
md4_ctx_t ctx = ctx0;
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);
@ -473,7 +473,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m18200_sxx (KERN_A
{
md4_ctx_t ctx = ctx0;
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
md4_update_global_utf16le (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len & 255);
md4_final (&ctx);

Loading…
Cancel
Save