github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-12-23 14:13:43 +01:00
Code Releases Activity

Fix SIMD issue

Browse Source
This commit is contained in:
Jens Steube 2017-07-10 12:05:37 +02:00
parent e70cc986da
commit e5a59a6611
1 changed files with 52 additions and 52 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

104
OpenCL/m09600.cl
View File

@ -243,22 +243,22 @@ __kernel void m09600_comp (__global pw_t *pws, __global const kernel_rule_t *rul
u32 w6[4]; u32 w6[4];
u32 w7[4]; u32 w7[4];
w0[0] = h32_from_64 (tmp[0]); w0[0] = h32_from_64_S (tmp[0]);
w0[1] = l32_from_64 (tmp[0]); w0[1] = l32_from_64_S (tmp[0]);
w0[2] = h32_from_64 (tmp[1]); w0[2] = h32_from_64_S (tmp[1]);
w0[3] = l32_from_64 (tmp[1]); w0[3] = l32_from_64_S (tmp[1]);
w1[0] = h32_from_64 (tmp[2]); w1[0] = h32_from_64_S (tmp[2]);
w1[1] = l32_from_64 (tmp[2]); w1[1] = l32_from_64_S (tmp[2]);
w1[2] = h32_from_64 (tmp[3]); w1[2] = h32_from_64_S (tmp[3]);
w1[3] = l32_from_64 (tmp[3]); w1[3] = l32_from_64_S (tmp[3]);
w2[0] = h32_from_64 (tmp[4]); w2[0] = h32_from_64_S (tmp[4]);
w2[1] = l32_from_64 (tmp[4]); w2[1] = l32_from_64_S (tmp[4]);
w2[2] = h32_from_64 (tmp[5]); w2[2] = h32_from_64_S (tmp[5]);
w2[3] = l32_from_64 (tmp[5]); w2[3] = l32_from_64_S (tmp[5]);
w3[0] = h32_from_64 (tmp[6]); w3[0] = h32_from_64_S (tmp[6]);
w3[1] = l32_from_64 (tmp[6]); w3[1] = l32_from_64_S (tmp[6]);
w3[2] = h32_from_64 (tmp[7]); w3[2] = h32_from_64_S (tmp[7]);
w3[3] = l32_from_64 (tmp[7]); w3[3] = l32_from_64_S (tmp[7]);
w4[0] = encryptedVerifierHashInputBlockKey[0]; w4[0] = encryptedVerifierHashInputBlockKey[0];
w4[1] = encryptedVerifierHashInputBlockKey[1]; w4[1] = encryptedVerifierHashInputBlockKey[1];
w4[2] = 0; w4[2] = 0;
@ -291,22 +291,22 @@ __kernel void m09600_comp (__global pw_t *pws, __global const kernel_rule_t *rul
digest0[2] = ctx.h[2]; digest0[2] = ctx.h[2];
digest0[3] = ctx.h[3]; digest0[3] = ctx.h[3];
w0[0] = h32_from_64 (tmp[0]); w0[0] = h32_from_64_S (tmp[0]);
w0[1] = l32_from_64 (tmp[0]); w0[1] = l32_from_64_S (tmp[0]);
w0[2] = h32_from_64 (tmp[1]); w0[2] = h32_from_64_S (tmp[1]);
w0[3] = l32_from_64 (tmp[1]); w0[3] = l32_from_64_S (tmp[1]);
w1[0] = h32_from_64 (tmp[2]); w1[0] = h32_from_64_S (tmp[2]);
w1[1] = l32_from_64 (tmp[2]); w1[1] = l32_from_64_S (tmp[2]);
w1[2] = h32_from_64 (tmp[3]); w1[2] = h32_from_64_S (tmp[3]);
w1[3] = l32_from_64 (tmp[3]); w1[3] = l32_from_64_S (tmp[3]);
w2[0] = h32_from_64 (tmp[4]); w2[0] = h32_from_64_S (tmp[4]);
w2[1] = l32_from_64 (tmp[4]); w2[1] = l32_from_64_S (tmp[4]);
w2[2] = h32_from_64 (tmp[5]); w2[2] = h32_from_64_S (tmp[5]);
w2[3] = l32_from_64 (tmp[5]); w2[3] = l32_from_64_S (tmp[5]);
w3[0] = h32_from_64 (tmp[6]); w3[0] = h32_from_64_S (tmp[6]);
w3[1] = l32_from_64 (tmp[6]); w3[1] = l32_from_64_S (tmp[6]);
w3[2] = h32_from_64 (tmp[7]); w3[2] = h32_from_64_S (tmp[7]);
w3[3] = l32_from_64 (tmp[7]); w3[3] = l32_from_64_S (tmp[7]);
w4[0] = encryptedVerifierHashValueBlockKey[0]; w4[0] = encryptedVerifierHashValueBlockKey[0];
w4[1] = encryptedVerifierHashValueBlockKey[1]; w4[1] = encryptedVerifierHashValueBlockKey[1];
w4[2] = 0; w4[2] = 0;
@ -341,14 +341,14 @@ __kernel void m09600_comp (__global pw_t *pws, __global const kernel_rule_t *rul
u32 ukey[8]; u32 ukey[8];
ukey[0] = h32_from_64 (digest0[0]); ukey[0] = h32_from_64_S (digest0[0]);
ukey[1] = l32_from_64 (digest0[0]); ukey[1] = l32_from_64_S (digest0[0]);
ukey[2] = h32_from_64 (digest0[1]); ukey[2] = h32_from_64_S (digest0[1]);
ukey[3] = l32_from_64 (digest0[1]); ukey[3] = l32_from_64_S (digest0[1]);
ukey[4] = h32_from_64 (digest0[2]); ukey[4] = h32_from_64_S (digest0[2]);
ukey[5] = l32_from_64 (digest0[2]); ukey[5] = l32_from_64_S (digest0[2]);
ukey[6] = h32_from_64 (digest0[3]); ukey[6] = h32_from_64_S (digest0[3]);
ukey[7] = l32_from_64 (digest0[3]); ukey[7] = l32_from_64_S (digest0[3]);
u32 ks[60]; u32 ks[60];
@ -420,21 +420,21 @@ __kernel void m09600_comp (__global pw_t *pws, __global const kernel_rule_t *rul
// encrypt with 2nd key // encrypt with 2nd key
ukey[0] = h32_from_64 (digest1[0]); ukey[0] = h32_from_64_S (digest1[0]);
ukey[1] = l32_from_64 (digest1[0]); ukey[1] = l32_from_64_S (digest1[0]);
ukey[2] = h32_from_64 (digest1[1]); ukey[2] = h32_from_64_S (digest1[1]);
ukey[3] = l32_from_64 (digest1[1]); ukey[3] = l32_from_64_S (digest1[1]);
ukey[4] = h32_from_64 (digest1[2]); ukey[4] = h32_from_64_S (digest1[2]);
ukey[5] = l32_from_64 (digest1[2]); ukey[5] = l32_from_64_S (digest1[2]);
ukey[6] = h32_from_64 (digest1[3]); ukey[6] = h32_from_64_S (digest1[3]);
ukey[7] = l32_from_64 (digest1[3]); ukey[7] = l32_from_64_S (digest1[3]);
AES256_set_encrypt_key (ks, ukey, s_te0, s_te1, s_te2, s_te3, s_te4); AES256_set_encrypt_key (ks, ukey, s_te0, s_te1, s_te2, s_te3, s_te4);
data[0] = h32_from_64 (digest[0]) ^ salt_bufs[salt_pos].salt_buf[0]; data[0] = h32_from_64_S (digest[0]) ^ salt_bufs[salt_pos].salt_buf[0];
data[1] = l32_from_64 (digest[0]) ^ salt_bufs[salt_pos].salt_buf[1]; data[1] = l32_from_64_S (digest[0]) ^ salt_bufs[salt_pos].salt_buf[1];
data[2] = h32_from_64 (digest[1]) ^ salt_bufs[salt_pos].salt_buf[2]; data[2] = h32_from_64_S (digest[1]) ^ salt_bufs[salt_pos].salt_buf[2];
data[3] = l32_from_64 (digest[1]) ^ salt_bufs[salt_pos].salt_buf[3]; data[3] = l32_from_64_S (digest[1]) ^ salt_bufs[salt_pos].salt_buf[3];
AES256_encrypt (ks, data, out, s_te0, s_te1, s_te2, s_te3, s_te4); AES256_encrypt (ks, data, out, s_te0, s_te1, s_te2, s_te3, s_te4);