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 old truncate_block()

Browse Source
This commit is contained in:
jsteube 2017-07-03 11:07:37 +02:00
parent 64704f36c2
commit b03382b334
7 changed files with 30 additions and 144 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

126
OpenCL/inc_common.cl
View File

@ -138,69 +138,6 @@ inline float get_entropy (const u32 *buf, const int elems)
* vector functions
*/
inline void truncate_block (u32x w[4], const u32 len)
{
switch (len)
{
case 0: w[0] &= 0;
w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 1: w[0] &= 0x000000FF;
w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 2: w[0] &= 0x0000FFFF;
w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 3: w[0] &= 0x00FFFFFF;
w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 4: w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 5: w[1] &= 0x000000FF;
w[2] &= 0;
w[3] &= 0;
break;
case 6: w[1] &= 0x0000FFFF;
w[2] &= 0;
w[3] &= 0;
break;
case 7: w[1] &= 0x00FFFFFF;
w[2] &= 0;
w[3] &= 0;
break;
case 8: w[2] &= 0;
w[3] &= 0;
break;
case 9: w[2] &= 0x000000FF;
w[3] &= 0;
break;
case 10: w[2] &= 0x0000FFFF;
w[3] &= 0;
break;
case 11: w[2] &= 0x00FFFFFF;
w[3] &= 0;
break;
case 12: w[3] &= 0;
break;
case 13: w[3] &= 0x000000FF;
break;
case 14: w[3] &= 0x0000FFFF;
break;
case 15: w[3] &= 0x00FFFFFF;
break;
}
}
inline void truncate_block_4x4_le (u32x w0[4], const u32 len)
{
switch (len)
@ -16235,69 +16172,6 @@ inline void append_0x80_8x4_S (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w
}
}
inline void truncate_block_S (u32 w[4], const u32 len)
{
switch (len)
{
case 0: w[0] &= 0;
w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 1: w[0] &= 0x000000FF;
w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 2: w[0] &= 0x0000FFFF;
w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 3: w[0] &= 0x00FFFFFF;
w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 4: w[1] &= 0;
w[2] &= 0;
w[3] &= 0;
break;
case 5: w[1] &= 0x000000FF;
w[2] &= 0;
w[3] &= 0;
break;
case 6: w[1] &= 0x0000FFFF;
w[2] &= 0;
w[3] &= 0;
break;
case 7: w[1] &= 0x00FFFFFF;
w[2] &= 0;
w[3] &= 0;
break;
case 8: w[2] &= 0;
w[3] &= 0;
break;
case 9: w[2] &= 0x000000FF;
w[3] &= 0;
break;
case 10: w[2] &= 0x0000FFFF;
w[3] &= 0;
break;
case 11: w[2] &= 0x00FFFFFF;
w[3] &= 0;
break;
case 12: w[3] &= 0;
break;
case 13: w[3] &= 0x000000FF;
break;
case 14: w[3] &= 0x0000FFFF;
break;
case 15: w[3] &= 0x00FFFFFF;
break;
}
}
inline void make_utf16be_S (const u32 in[4], u32 out1[4], u32 out2[4])
{
#ifdef IS_NV

2
OpenCL/m06800.cl
View File

@ -1556,7 +1556,7 @@ __kernel void m06800_comp (__global pw_t *pws, __global const kernel_rule_t *rul
out[2] = swap32 (out[2]);
out[3] = swap32 (out[3]);
truncate_block (out, salt_len);
truncate_block_4x4_le (out, salt_len);
if ((out[0] == salt_buf[0])
&& (out[1] == salt_buf[1])

6
OpenCL/m07400.cl
View File

@ -868,7 +868,7 @@ __kernel void m07400_init (__global pw_t *pws, __global const kernel_rule_t *rul
alt_result_tmp[6] = 0;
alt_result_tmp[7] = 0;
truncate_block (alt_result_tmp, pw_len);
truncate_block_4x4_le (alt_result_tmp, pw_len);
/* Add the key string. */
@ -978,7 +978,7 @@ __kernel void m07400_init (__global pw_t *pws, __global const kernel_rule_t *rul
bswap8 (p_bytes);
truncate_block (p_bytes, pw_len);
truncate_block_4x4_le (p_bytes, pw_len);
tmps[gid].p_bytes[0] = p_bytes[0];
tmps[gid].p_bytes[1] = p_bytes[1];
@ -1021,7 +1021,7 @@ __kernel void m07400_init (__global pw_t *pws, __global const kernel_rule_t *rul
bswap8 (s_bytes);
truncate_block (s_bytes, salt_len);
truncate_block_4x4_le (s_bytes, salt_len);
tmps[gid].s_bytes[0] = s_bytes[0];
tmps[gid].s_bytes[1] = s_bytes[1];

12
OpenCL/m13100_a0.cl
View File

@ -481,7 +481,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block (w0, edata2_left & 0xf);
truncate_block_4x4_le (w0, edata2_left & 0xf);
append_0x80_1x4 (w0, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -494,7 +495,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block (w1, edata2_left & 0xf);
truncate_block_4x4_le (w1, edata2_left & 0xf);
append_0x80_1x4 (w1, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -508,7 +510,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block (w2, edata2_left & 0xf);
truncate_block_4x4_le (w2, edata2_left & 0xf);
append_0x80_1x4 (w2, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -523,7 +526,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block (w3, edata2_left & 0xf);
truncate_block_4x4_le (w3, edata2_left & 0xf);
append_0x80_1x4 (w3, edata2_left & 0xf);
if (edata2_left < 56)

12
OpenCL/m13100_a1.cl
View File

@ -479,7 +479,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block (w0, edata2_left & 0xf);
truncate_block_4x4_le (w0, edata2_left & 0xf);
append_0x80_1x4 (w0, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -492,7 +493,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block (w1, edata2_left & 0xf);
truncate_block_4x4_le (w1, edata2_left & 0xf);
append_0x80_1x4 (w1, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -506,7 +508,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block (w2, edata2_left & 0xf);
truncate_block_4x4_le (w2, edata2_left & 0xf);
append_0x80_1x4 (w2, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -521,7 +524,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block (w3, edata2_left & 0xf);
truncate_block_4x4_le (w3, edata2_left & 0xf);
append_0x80_1x4 (w3, edata2_left & 0xf);
if (edata2_left < 56)

12
OpenCL/m13100_a3.cl
View File

@ -479,7 +479,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block (w0, edata2_left & 0xf);
truncate_block_4x4_le (w0, edata2_left & 0xf);
append_0x80_1x4 (w0, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -492,7 +493,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block (w1, edata2_left & 0xf);
truncate_block_4x4_le (w1, edata2_left & 0xf);
append_0x80_1x4 (w1, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -506,7 +508,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block (w2, edata2_left & 0xf);
truncate_block_4x4_le (w2, edata2_left & 0xf);
append_0x80_1x4 (w2, edata2_left & 0xf);
w3[2] = (64 + edata2_len) * 8;
@ -521,7 +524,8 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block (w3, edata2_left & 0xf);
truncate_block_4x4_le (w3, edata2_left & 0xf);
append_0x80_1x4 (w3, edata2_left & 0xf);
if (edata2_left < 56)

4
OpenCL/m13400.cl
View File

@ -939,7 +939,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul
// we need to clear the buffer of the padding data
truncate_block (out, 16 - pad_byte);
truncate_block_4x4_le (out, 16 - pad_byte);
// it's also a good point to push our 0x80
@ -1146,7 +1146,7 @@ __kernel void m13400_comp (__global pw_t *pws, __global const kernel_rule_t *rul
// we need to clear the buffer of the padding data
truncate_block (out, 16 - pad_byte);
truncate_block_4x4_le (out, 16 - pad_byte);
// it's also a good point to push our 0x80