arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-26 01:50:10 +00:00
Code Issues Releases Wiki Activity

Use truncate_block_4x4_le_S() instead of truncate_block_4x4_le() in -m 6800

Browse Source
This commit is contained in:
jsteube 2017-08-24 19:53:29 +02:00
parent 9b73c464d2
commit 51dc1c7db3
2 changed files with 111 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

110
OpenCL/inc_common.cl
View File

@ -31506,6 +31506,116 @@ static void switch_buffer_by_offset_1x64_be (u32x w[64], const u32 offset)
* vector functions as scalar (for outer loop usage)
*/
static void truncate_block_4x4_le_S (u32 w0[4], const u32 len)
{
switch (len)
{
case 0:
w0[0] = 0;
w0[1] = 0;
w0[2] = 0;
w0[3] = 0;
break;
case 1:
w0[0] &= 0x000000ff;
w0[1] = 0;
w0[2] = 0;
w0[3] = 0;
break;
case 2:
w0[0] &= 0x0000ffff;
w0[1] = 0;
w0[2] = 0;
w0[3] = 0;
break;
case 3:
w0[0] &= 0x00ffffff;
w0[1] = 0;
w0[2] = 0;
w0[3] = 0;
break;
case 4:
w0[1] = 0;
w0[2] = 0;
w0[3] = 0;
break;
case 5:
w0[1] &= 0x000000ff;
w0[2] = 0;
w0[3] = 0;
break;
case 6:
w0[1] &= 0x0000ffff;
w0[2] = 0;
w0[3] = 0;
break;
case 7:
w0[1] &= 0x00ffffff;
w0[2] = 0;
w0[3] = 0;
break;
case 8:
w0[2] = 0;
w0[3] = 0;
break;
case 9:
w0[2] &= 0x000000ff;
w0[3] = 0;
break;
case 10:
w0[2] &= 0x0000ffff;
w0[3] = 0;
break;
case 11:
w0[2] &= 0x00ffffff;
w0[3] = 0;
break;
case 12:
w0[3] = 0;
break;
case 13:
w0[3] &= 0x000000ff;
break;
case 14:
w0[3] &= 0x0000ffff;
break;
case 15:
w0[3] &= 0x00ffffff;
break;
}
}
static void append_0x01_2x4_S (u32 w0[4], u32 w1[4], const u32 offset)
{
const u32 tmp = 0x01 << ((offset & 3) * 8);

2
OpenCL/m06800.cl
View File

@ -356,7 +356,7 @@ __kernel void m06800_comp (__global pw_t *pws, __constant const kernel_rule_t *r
out[2] = swap32_S (out[2]);
out[3] = swap32_S (out[3]);
truncate_block_4x4_le (out, salt_len);
truncate_block_4x4_le_S (out, salt_len);
if ((out[0] == salt_buf[0])
&& (out[1] == salt_buf[1])