mirror of
https://github.com/hashcat/hashcat.git
synced 2024-12-22 22:58:30 +00:00
Give the compiler a hint for automatic optimizations based on password length
This commit is contained in:
parent
c9da60c73a
commit
0e428b3c40
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&md5_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
md5_final (&md5_ctx);
|
||||
|
||||
|
@ -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);
|
||||
md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md4_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha224_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha512_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md4_update_global_utf16le (&md4_ctx1, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
md4_final (&md4_ctx1);
|
||||
|
||||
|
@ -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);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -43,7 +43,7 @@ __kernel void m04900_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_update (&ctx0, s, salt_len);
|
||||
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -110,7 +110,7 @@ __kernel void m04900_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_update (&ctx0, s, salt_len);
|
||||
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -32,7 +32,7 @@ __kernel void m05100_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -90,7 +90,7 @@ __kernel void m05100_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -30,7 +30,7 @@ __kernel void m05200_init (KERN_ATTR_TMPS (pwsafe3_tmp_t))
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha256_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
|
@ -545,7 +545,7 @@ __kernel void m05500_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
md4_init (&ctx0);
|
||||
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -679,7 +679,7 @@ __kernel void m05500_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
md4_init (&ctx0);
|
||||
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -33,7 +33,7 @@ __kernel void m05600_mxx (KERN_ATTR_ESALT (netntlm_t))
|
||||
|
||||
md4_init (&ctx10);
|
||||
|
||||
md4_update_global_utf16le (&ctx10, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&ctx10, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -142,7 +142,7 @@ __kernel void m05600_sxx (KERN_ATTR_ESALT (netntlm_t))
|
||||
|
||||
md4_init (&ctx10);
|
||||
|
||||
md4_update_global_utf16le (&ctx10, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&ctx10, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -2214,7 +2214,7 @@ __kernel void m05800_init (KERN_ATTR_TMPS (androidpin_tmp_t))
|
||||
|
||||
sha1_update_64 (&ctx, w0, w1, w2, w3, 1);
|
||||
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha1_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
|
@ -32,7 +32,7 @@ __kernel void m06000_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
ripemd160_init (&ctx0);
|
||||
|
||||
ripemd160_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
ripemd160_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -86,7 +86,7 @@ __kernel void m06000_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
ripemd160_init (&ctx0);
|
||||
|
||||
ripemd160_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
ripemd160_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -63,7 +63,7 @@ __kernel void m06100_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
whirlpool_init (&ctx0, s_Ch, s_Cl);
|
||||
|
||||
whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -148,7 +148,7 @@ __kernel void m06100_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
whirlpool_init (&ctx0, s_Ch, s_Cl);
|
||||
|
||||
whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
whirlpool_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -70,7 +70,7 @@ __kernel void m06400_init (KERN_ATTR_TMPS (sha256aix_tmp_t))
|
||||
|
||||
sha256_hmac_ctx_t sha256_hmac_ctx;
|
||||
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
|
||||
|
@ -86,7 +86,7 @@ __kernel void m06500_init (KERN_ATTR_TMPS (sha512aix_tmp_t))
|
||||
|
||||
sha512_hmac_ctx_t sha512_hmac_ctx;
|
||||
|
||||
sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
|
||||
|
@ -65,7 +65,7 @@ __kernel void m06600_init (KERN_ATTR_TMPS (agilekey_tmp_t))
|
||||
|
||||
sha1_hmac_ctx_t sha1_hmac_ctx;
|
||||
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
@ -64,7 +64,7 @@ __kernel void m06700_init (KERN_ATTR_TMPS (sha1aix_tmp_t))
|
||||
|
||||
sha1_hmac_ctx_t sha1_hmac_ctx;
|
||||
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
@ -71,7 +71,7 @@ __kernel void m06800_init (KERN_ATTR_TMPS (lastpass_tmp_t))
|
||||
|
||||
sha256_hmac_ctx_t sha256_hmac_ctx;
|
||||
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
|
||||
|
@ -34,7 +34,7 @@ __kernel void m07000_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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -118,7 +118,7 @@ __kernel void m07000_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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -86,7 +86,7 @@ __kernel void m07100_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, pbkdf2_sha
|
||||
|
||||
sha512_hmac_ctx_t sha512_hmac_ctx;
|
||||
|
||||
sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
|
||||
|
@ -307,7 +307,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_mxx (KERN_A
|
||||
|
||||
md4_init (&ctx0);
|
||||
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -376,7 +376,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m07500_sxx (KERN_A
|
||||
|
||||
md4_init (&ctx0);
|
||||
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -29,7 +29,7 @@ __kernel void m07900_init (KERN_ATTR_TMPS (drupal7_tmp_t))
|
||||
|
||||
sha512_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
|
@ -34,7 +34,7 @@ __kernel void m08100_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_update_global (&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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -90,7 +90,7 @@ __kernel void m08100_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_update_global (&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);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -87,7 +87,7 @@ __kernel void m08200_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha512_tmp_t, cloudkey_t
|
||||
|
||||
sha512_hmac_ctx_t sha512_hmac_ctx;
|
||||
|
||||
sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
|
||||
|
@ -62,7 +62,7 @@ __kernel void m08300_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
ctx1.len = 1;
|
||||
|
||||
sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len);
|
||||
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);
|
||||
|
||||
@ -177,7 +177,7 @@ __kernel void m08300_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
ctx1.len = 1;
|
||||
|
||||
sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len);
|
||||
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);
|
||||
|
||||
|
@ -68,7 +68,7 @@ __kernel void m08400_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx1l);
|
||||
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -229,7 +229,7 @@ __kernel void m08400_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx1l);
|
||||
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -63,7 +63,7 @@ __kernel void m08800_init (KERN_ATTR_TMPS_ESALT (androidfde_tmp_t, androidfde_t)
|
||||
|
||||
sha1_hmac_ctx_t sha1_hmac_ctx;
|
||||
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
@ -219,7 +219,7 @@ __kernel void m08900_init (KERN_ATTR_TMPS (scrypt_tmp_t))
|
||||
|
||||
sha256_hmac_ctx_t sha256_hmac_ctx;
|
||||
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha256_hmac_update_global_swap (&sha256_hmac_ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
@ -327,7 +327,7 @@ __kernel void m08900_comp (KERN_ATTR_TMPS (scrypt_tmp_t))
|
||||
|
||||
sha256_hmac_ctx_t ctx;
|
||||
|
||||
sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_hmac_init_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
for (u32 l = 0; l < SCRYPT_CNT4; l += 4)
|
||||
{
|
||||
|
@ -33,7 +33,7 @@ __kernel void m09400_init (KERN_ATTR_TMPS_ESALT (office2007_tmp_t, office2007_t)
|
||||
|
||||
sha1_update_global (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
@ -33,7 +33,7 @@ __kernel void m09500_init (KERN_ATTR_TMPS_ESALT (office2010_tmp_t, office2010_t)
|
||||
|
||||
sha1_update_global (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
@ -33,7 +33,7 @@ __kernel void m09600_init (KERN_ATTR_TMPS_ESALT (office2013_tmp_t, office2013_t)
|
||||
|
||||
sha512_update_global (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
sha512_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha512_update_global_utf16le_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
|
@ -32,7 +32,7 @@ __kernel void m09900_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -86,7 +86,7 @@ __kernel void m09900_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
md5_init (&ctx0);
|
||||
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -27,7 +27,7 @@ __kernel void m10300_init (KERN_ATTR_TMPS (saph_sha1_tmp_t))
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha1_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
@ -58,7 +58,7 @@ __kernel void m10300_loop (KERN_ATTR_TMPS (saph_sha1_tmp_t))
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* load
|
||||
|
@ -529,7 +529,7 @@ __kernel void m10700_init (KERN_ATTR_TMPS_ESALT (pdf17l8_tmp_t, pdf_t))
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha256_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
|
@ -1118,7 +1118,7 @@ __kernel void m10700_init (KERN_ATTR_TMPS_ESALT (pdf17l8_tmp_t, pdf_t))
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha256_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
|
@ -32,7 +32,7 @@ __kernel void m10800_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha384_init (&ctx0);
|
||||
|
||||
sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -86,7 +86,7 @@ __kernel void m10800_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha384_init (&ctx0);
|
||||
|
||||
sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha384_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -70,7 +70,7 @@ __kernel void m10900_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha
|
||||
|
||||
sha256_hmac_ctx_t sha256_hmac_ctx;
|
||||
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
|
||||
|
@ -34,7 +34,7 @@ __kernel void m11000_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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -90,7 +90,7 @@ __kernel void m11000_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);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -84,7 +84,7 @@ __kernel void m11100_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
md5_init (&ctx0t);
|
||||
|
||||
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len);
|
||||
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -240,7 +240,7 @@ __kernel void m11100_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
md5_init (&ctx0t);
|
||||
|
||||
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len);
|
||||
md5_update_global (&ctx0t, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -38,7 +38,7 @@ __kernel void m11200_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx2l);
|
||||
|
||||
sha1_update_global_swap (&ctx2l, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx2l, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -164,7 +164,7 @@ __kernel void m11200_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx2l);
|
||||
|
||||
sha1_update_global_swap (&ctx2l, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx2l, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -86,7 +86,7 @@ __kernel void m11300_init (KERN_ATTR_TMPS_ESALT (bitcoin_wallet_tmp_t, bitcoin_w
|
||||
|
||||
sha512_init (&ctx);
|
||||
|
||||
sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha512_update_global_swap (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
|
@ -64,7 +64,7 @@ __kernel void m11400_mxx (KERN_ATTR_ESALT (sip_t))
|
||||
|
||||
md5_update_global (&ctx0, esalt_bufs[digests_offset].salt_buf, esalt_bufs[digests_offset].salt_len);
|
||||
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -170,7 +170,7 @@ __kernel void m11400_sxx (KERN_ATTR_ESALT (sip_t))
|
||||
|
||||
md5_update_global (&ctx0, esalt_bufs[digests_offset].salt_buf, esalt_bufs[digests_offset].salt_len);
|
||||
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md5_update_global (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -61,7 +61,7 @@ __kernel void m11700_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
streebog256_init (&ctx0, s_sbob_sl64);
|
||||
|
||||
streebog256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
streebog256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -144,7 +144,7 @@ __kernel void m11700_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
streebog256_init (&ctx0, s_sbob_sl64);
|
||||
|
||||
streebog256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
streebog256_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -61,7 +61,7 @@ __kernel void m11800_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
streebog512_init (&ctx0, s_sbob_sl64);
|
||||
|
||||
streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -144,7 +144,7 @@ __kernel void m11800_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
streebog512_init (&ctx0, s_sbob_sl64);
|
||||
|
||||
streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
streebog512_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -62,7 +62,7 @@ __kernel void m11900_init (KERN_ATTR_TMPS_ESALT (pbkdf2_md5_tmp_t, pbkdf2_md5_t)
|
||||
|
||||
md5_hmac_ctx_t md5_hmac_ctx;
|
||||
|
||||
md5_hmac_init_global (&md5_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
md5_hmac_init_global (&md5_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = md5_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = md5_hmac_ctx.ipad.h[1];
|
||||
|
@ -64,7 +64,7 @@ __kernel void m12000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, pbkdf2_sha1_
|
||||
|
||||
sha1_hmac_ctx_t sha1_hmac_ctx;
|
||||
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
@ -32,7 +32,7 @@ __kernel void m12200_init (KERN_ATTR_TMPS (ecryptfs_tmp_t))
|
||||
|
||||
sha512_update_global (&ctx, salt_bufs[salt_pos].salt_buf, salt_bufs[salt_pos].salt_len);
|
||||
|
||||
sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha512_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha512_final (&ctx);
|
||||
|
||||
|
@ -86,7 +86,7 @@ __kernel void m12300_init (KERN_ATTR_TMPS (oraclet_tmp_t))
|
||||
|
||||
sha512_hmac_ctx_t sha512_hmac_ctx;
|
||||
|
||||
sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha512_hmac_init_global_swap (&sha512_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha512_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha512_hmac_ctx.ipad.h[1];
|
||||
|
@ -78,7 +78,7 @@ __kernel void m12600_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -224,7 +224,7 @@ __kernel void m12600_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -65,7 +65,7 @@ __kernel void m12700_init (KERN_ATTR_TMPS (mywallet_tmp_t))
|
||||
|
||||
sha1_hmac_ctx_t sha1_hmac_ctx;
|
||||
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
@ -96,7 +96,7 @@ __kernel void m12800_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha
|
||||
|
||||
md4_init (&md4_ctx);
|
||||
|
||||
md4_update_global_utf16le (&md4_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&md4_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
md4_final (&md4_ctx);
|
||||
|
||||
|
@ -70,7 +70,7 @@ __kernel void m12900_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha
|
||||
|
||||
sha256_hmac_ctx_t sha256_hmac_ctx;
|
||||
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
|
||||
|
@ -70,7 +70,7 @@ __kernel void m13000_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha256_tmp_t, pbkdf2_sha
|
||||
|
||||
sha256_hmac_ctx_t sha256_hmac_ctx;
|
||||
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_hmac_init_global_swap (&sha256_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha256_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha256_hmac_ctx.ipad.h[1];
|
||||
|
@ -405,7 +405,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_mxx (KERN_A
|
||||
|
||||
md4_init (&ctx0);
|
||||
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -465,7 +465,7 @@ __kernel void __attribute__((reqd_work_group_size(64, 1, 1))) m13100_sxx (KERN_A
|
||||
|
||||
md4_init (&ctx0);
|
||||
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
md4_update_global_utf16le (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -29,7 +29,7 @@ __kernel void m13200_init (KERN_ATTR_TMPS (axcrypt_tmp_t))
|
||||
|
||||
sha1_init (&ctx);
|
||||
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha1_final (&ctx);
|
||||
|
||||
|
@ -32,7 +32,7 @@ __kernel void m13300_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -88,7 +88,7 @@ __kernel void m13300_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx0);
|
||||
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -29,7 +29,7 @@ __kernel void m13400_init (KERN_ATTR_TMPS_ESALT (keepass_tmp_t, keepass_t))
|
||||
|
||||
sha256_init (&ctx);
|
||||
|
||||
sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
sha256_final (&ctx);
|
||||
|
||||
|
@ -61,7 +61,7 @@ __kernel void m13500_mxx (KERN_ATTR_ESALT (pstoken_t))
|
||||
* base
|
||||
*/
|
||||
|
||||
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -144,7 +144,7 @@ __kernel void m13500_sxx (KERN_ATTR_ESALT (pstoken_t))
|
||||
* base
|
||||
*/
|
||||
|
||||
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -64,7 +64,7 @@ __kernel void m13600_init (KERN_ATTR_TMPS_ESALT (pbkdf2_sha1_tmp_t, zip2_t))
|
||||
|
||||
sha1_hmac_ctx_t sha1_hmac_ctx;
|
||||
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
@ -32,7 +32,7 @@ __kernel void m13800_mxx (KERN_ATTR_ESALT (win8phone_t))
|
||||
|
||||
sha256_init (&ctx0);
|
||||
|
||||
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -88,7 +88,7 @@ __kernel void m13800_sxx (KERN_ATTR_ESALT (win8phone_t))
|
||||
|
||||
sha256_init (&ctx0);
|
||||
|
||||
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len);
|
||||
sha256_update_global_utf16le_swap (&ctx0, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -68,7 +68,7 @@ __kernel void m13900_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx1l);
|
||||
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
@ -229,7 +229,7 @@ __kernel void m13900_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_init (&ctx1l);
|
||||
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_update_global_swap (&ctx1l, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
/**
|
||||
* loop
|
||||
|
@ -139,7 +139,7 @@ __kernel void m14400_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_update_64 (&ctx1, d20, d21, d22, d23, 2);
|
||||
|
||||
sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len);
|
||||
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);
|
||||
|
||||
@ -229,7 +229,7 @@ __kernel void m14400_mxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_update_64 (&ctx, d20, d21, d22, d23, 2);
|
||||
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
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);
|
||||
|
||||
@ -396,7 +396,7 @@ __kernel void m14400_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_update_64 (&ctx1, d20, d21, d22, d23, 2);
|
||||
|
||||
sha1_update_global_swap (&ctx1, pws[gid].i, pws[gid].pw_len);
|
||||
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);
|
||||
|
||||
@ -486,7 +486,7 @@ __kernel void m14400_sxx (KERN_ATTR_BASIC ())
|
||||
|
||||
sha1_update_64 (&ctx, d20, d21, d22, d23, 2);
|
||||
|
||||
sha1_update_global_swap (&ctx, pws[gid].i, pws[gid].pw_len);
|
||||
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);
|
||||
|
||||
|
@ -73,7 +73,7 @@ __kernel void m14611_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t))
|
||||
|
||||
sha1_hmac_ctx_t sha1_hmac_ctx;
|
||||
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad32[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad32[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
@ -73,7 +73,7 @@ __kernel void m14612_init (KERN_ATTR_TMPS_ESALT (luks_tmp_t, luks_t))
|
||||
|
||||
sha1_hmac_ctx_t sha1_hmac_ctx;
|
||||
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len);
|
||||
sha1_hmac_init_global_swap (&sha1_hmac_ctx, pws[gid].i, pws[gid].pw_len & 255);
|
||||
|
||||
tmps[gid].ipad32[0] = sha1_hmac_ctx.ipad.h[0];
|
||||
tmps[gid].ipad32[1] = sha1_hmac_ctx.ipad.h[1];
|
||||
|
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue
Block a user