github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-11-17 08:28:39 +01:00
Code Releases Activity

Get rid of swap32() in make_kn() in -m 2500

Browse Source
This commit is contained in:
jsteube 2017-09-19 13:45:43 +02:00
parent ab1dabebbe
commit 68f5b12754
2 changed files with 68 additions and 104 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

86
OpenCL/m02500.cl
View File

@ -19,6 +19,36 @@
#define COMPARE_S "inc_comp_single.cl" #define COMPARE_S "inc_comp_single.cl"
#define COMPARE_M "inc_comp_multi.cl" #define COMPARE_M "inc_comp_multi.cl"
void make_kn (u32 *k)
{
u32 kl[4];
u32 kr[4];
kl[0] = (k[0] << 1) & 0xfefefefe;
kl[1] = (k[1] << 1) & 0xfefefefe;
kl[2] = (k[2] << 1) & 0xfefefefe;
kl[3] = (k[3] << 1) & 0xfefefefe;
kr[0] = (k[0] >> 7) & 0x01010101;
kr[1] = (k[1] >> 7) & 0x01010101;
kr[2] = (k[2] >> 7) & 0x01010101;
kr[3] = (k[3] >> 7) & 0x01010101;
const u32 c = kr[0] & 1;
kr[0] = kr[0] >> 8 | kr[1] << 24;
kr[1] = kr[1] >> 8 | kr[2] << 24;
kr[2] = kr[2] >> 8 | kr[3] << 24;
kr[3] = kr[3] >> 8;
kr[3] ^= c * 0x87000000;
k[0] = kl[0] | kr[0];
k[1] = kl[1] | kr[1];
k[2] = kl[2] | kr[2];
k[3] = kl[3] | kr[3];
}
void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5], u32x digest[5]) void hmac_sha1_run_V (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[5], u32x opad[5], u32x digest[5])
{ {
digest[0] = ipad[0]; digest[0] = ipad[0];
@ -557,37 +587,13 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul
aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4);
k[0] = swap32_S (k[0]); make_kn (k);
k[1] = swap32_S (k[1]);
k[2] = swap32_S (k[2]);
k[3] = swap32_S (k[3]);
const u32 c1 = k[3] >> 31;
k[3] = (k[3] << 1) | (k[2] >> 31);
k[2] = (k[2] << 1) | (k[1] >> 31);
k[1] = (k[1] << 1) | (k[0] >> 31);
k[0] = (k[0] << 1);
k[0] ^= c1 * 0x87;
if (eapol_left < 16) if (eapol_left < 16)
{ {
const u32 c2 = k[3] >> 31; make_kn (k);
k[3] = (k[3] << 1) | (k[2] >> 31);
k[2] = (k[2] << 1) | (k[1] >> 31);
k[1] = (k[1] << 1) | (k[0] >> 31);
k[0] = (k[0] << 1);
k[0] ^= c2 * 0x87;
} }
k[0] = swap32_S (k[0]);
k[1] = swap32_S (k[1]);
k[2] = swap32_S (k[2]);
k[3] = swap32_S (k[3]);
m[0] ^= k[0]; m[0] ^= k[0];
m[1] ^= k[1]; m[1] ^= k[1];
m[2] ^= k[2]; m[2] ^= k[2];
@ -836,37 +842,13 @@ __kernel void m02500_comp (__global pw_t *pws, __global const kernel_rule_t *rul
aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4);
k[0] = swap32_S (k[0]); make_kn (k);
k[1] = swap32_S (k[1]);
k[2] = swap32_S (k[2]);
k[3] = swap32_S (k[3]);
const u32 c1 = k[3] >> 31;
k[3] = (k[3] << 1) | (k[2] >> 31);
k[2] = (k[2] << 1) | (k[1] >> 31);
k[1] = (k[1] << 1) | (k[0] >> 31);
k[0] = (k[0] << 1);
k[0] ^= c1 * 0x87;
if (eapol_left < 16) if (eapol_left < 16)
{ {
const u32 c2 = k[3] >> 31; make_kn (k);
k[3] = (k[3] << 1) | (k[2] >> 31);
k[2] = (k[2] << 1) | (k[1] >> 31);
k[1] = (k[1] << 1) | (k[0] >> 31);
k[0] = (k[0] << 1);
k[0] ^= c2 * 0x87;
} }
k[0] = swap32_S (k[0]);
k[1] = swap32_S (k[1]);
k[2] = swap32_S (k[2]);
k[3] = swap32_S (k[3]);
m[0] ^= k[0]; m[0] ^= k[0];
m[1] ^= k[1]; m[1] ^= k[1];
m[2] ^= k[2]; m[2] ^= k[2];

86
OpenCL/m02501.cl
View File

@ -34,6 +34,36 @@ u8 hex_to_u8 (const u8 hex[2])
return (v); return (v);
} }
void make_kn (u32 *k)
{
u32 kl[4];
u32 kr[4];
kl[0] = (k[0] << 1) & 0xfefefefe;
kl[1] = (k[1] << 1) & 0xfefefefe;
kl[2] = (k[2] << 1) & 0xfefefefe;
kl[3] = (k[3] << 1) & 0xfefefefe;
kr[0] = (k[0] >> 7) & 0x01010101;
kr[1] = (k[1] >> 7) & 0x01010101;
kr[2] = (k[2] >> 7) & 0x01010101;
kr[3] = (k[3] >> 7) & 0x01010101;
const u32 c = kr[0] & 1;
kr[0] = kr[0] >> 8 | kr[1] << 24;
kr[1] = kr[1] >> 8 | kr[2] << 24;
kr[2] = kr[2] >> 8 | kr[3] << 24;
kr[3] = kr[3] >> 8;
kr[3] ^= c * 0x87000000;
k[0] = kl[0] | kr[0];
k[1] = kl[1] | kr[1];
k[2] = kl[2] | kr[2];
k[3] = kl[3] | kr[3];
}
__kernel void m02501_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const wpa_t *wpa_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max) __kernel void m02501_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global wpapmk_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global const wpa_t *wpa_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u64 gid_max)
{ {
const u64 gid = get_global_id (0); const u64 gid = get_global_id (0);
@ -432,37 +462,13 @@ __kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rul
aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4);
k[0] = swap32_S (k[0]); make_kn (k);
k[1] = swap32_S (k[1]);
k[2] = swap32_S (k[2]);
k[3] = swap32_S (k[3]);
const u32 c1 = k[3] >> 31;
k[3] = (k[3] << 1) | (k[2] >> 31);
k[2] = (k[2] << 1) | (k[1] >> 31);
k[1] = (k[1] << 1) | (k[0] >> 31);
k[0] = (k[0] << 1);
k[0] ^= c1 * 0x87;
if (eapol_left < 16) if (eapol_left < 16)
{ {
const u32 c2 = k[3] >> 31; make_kn (k);
k[3] = (k[3] << 1) | (k[2] >> 31);
k[2] = (k[2] << 1) | (k[1] >> 31);
k[1] = (k[1] << 1) | (k[0] >> 31);
k[0] = (k[0] << 1);
k[0] ^= c2 * 0x87;
} }
k[0] = swap32_S (k[0]);
k[1] = swap32_S (k[1]);
k[2] = swap32_S (k[2]);
k[3] = swap32_S (k[3]);
m[0] ^= k[0]; m[0] ^= k[0];
m[1] ^= k[1]; m[1] ^= k[1];
m[2] ^= k[2]; m[2] ^= k[2];
@ -711,37 +717,13 @@ __kernel void m02501_comp (__global pw_t *pws, __global const kernel_rule_t *rul
aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4); aes128_encrypt (ks, k, k, s_te0, s_te1, s_te2, s_te3, s_te4);
k[0] = swap32_S (k[0]); make_kn (k);
k[1] = swap32_S (k[1]);
k[2] = swap32_S (k[2]);
k[3] = swap32_S (k[3]);
const u32 c1 = k[3] >> 31;
k[3] = (k[3] << 1) | (k[2] >> 31);
k[2] = (k[2] << 1) | (k[1] >> 31);
k[1] = (k[1] << 1) | (k[0] >> 31);
k[0] = (k[0] << 1);
k[0] ^= c1 * 0x87;
if (eapol_left < 16) if (eapol_left < 16)
{ {
const u32 c2 = k[3] >> 31; make_kn (k);
k[3] = (k[3] << 1) | (k[2] >> 31);
k[2] = (k[2] << 1) | (k[1] >> 31);
k[1] = (k[1] << 1) | (k[0] >> 31);
k[0] = (k[0] << 1);
k[0] ^= c2 * 0x87;
} }
k[0] = swap32_S (k[0]);
k[1] = swap32_S (k[1]);
k[2] = swap32_S (k[2]);
k[3] = swap32_S (k[3]);
m[0] ^= k[0]; m[0] ^= k[0];
m[1] ^= k[1]; m[1] ^= k[1];
m[2] ^= k[2]; m[2] ^= k[2];