arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-22 22:58:30 +00:00
Code Issues Releases Wiki Activity

Add different code generators for truncate_block(), add results to inc_common.cl and make use of them in m01800-pure.cl

Browse Source
This commit is contained in:
jsteube 2017-07-03 10:41:09 +02:00
parent b9b2112b64
commit a22da36a00
6 changed files with 2000 additions and 46 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1843
OpenCL/inc_common.cl
View File

File diff suppressed because it is too large Load Diff

35
OpenCL/m01800-pure.cl
View File

@ -118,14 +118,9 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul
#ifdef _unroll
#pragma unroll
#endif
for (int i = 0; i < 16; i++) t_final[i] = swap32 (final[i]);
for (int i = 0; i < 16; i++) t_final[i] = final[i];
truncate_block_64 (t_final, pl);
#ifdef _unroll
#pragma unroll
#endif
for (int i = 0; i < 16; i++) t_final[i] = swap32 (t_final[i]);
truncate_block_16x4_be (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
sha512_update (&ctx, t_final, pl);
@ -212,17 +207,7 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul
p_final[idx + 15] = final[15];
}
#ifdef _unroll
#pragma unroll
#endif
for (int i = 0; i < 16; i++) final[i] = swap32 (final[i]);
truncate_block_64 (final, pl);
#ifdef _unroll
#pragma unroll
#endif
for (int i = 0; i < 16; i++) final[i] = swap32 (final[i]);
truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
p_final[idx + 0] = final[ 0];
p_final[idx + 1] = final[ 1];
@ -250,7 +235,7 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul
sha512_init (&ctx);
for (u32 j = 0; j < 16 + ((tmps[gid].alt_result[ 0] >> 24) & 0xff); j++)
for (u32 j = 0; j < 16 + (tmps[gid].alt_result[0] >> 24); j++)
{
sha512_update (&ctx, s, salt_len);
}
@ -296,17 +281,7 @@ __kernel void m01800_init (__global pw_t *pws, __global const kernel_rule_t *rul
s_final[idx + 15] = final[15];
}
#ifdef _unroll
#pragma unroll
#endif
for (int i = 0; i < 16; i++) final[i] = swap32 (final[i]);
truncate_block_64 (final, pl);
#ifdef _unroll
#pragma unroll
#endif
for (int i = 0; i < 16; i++) final[i] = swap32 (final[i]);
truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
s_final[idx + 0] = final[ 0];
s_final[idx + 1] = final[ 1];

42
tools/code_generators/GEN_truncate_block_16x4_be.pl Normal file
View File

@ -0,0 +1,42 @@
#!/usr/bin/perl
use strict;
use warnings;
for (my $i = 0; $i < 64; $i++)
{
printf (" case %2d:\n", $i);
my $id4 = int ($i / 4);
my $im4 = int ($i % 4);
if ($im4 == 0)
{
printf (" w%d[%d] = 0;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 1)
{
printf (" w%d[%d] &= 0xff000000;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 2)
{
printf (" w%d[%d] &= 0xffff0000;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 3)
{
printf (" w%d[%d] &= 0xffffff00;\n", $id4 / 4, $id4 % 4);
}
for (my $j = $id4 + 1; $j < 16; $j++)
{
my $jd4 = int ($j / 4);
my $jm4 = int ($j % 4);
printf (" w%d[%d] = 0;\n", $jd4, $jm4);
}
printf ("\n");
printf (" break;\n");
printf ("\n");
}

42
tools/code_generators/GEN_truncate_block_16x4_le.pl Normal file
View File

@ -0,0 +1,42 @@
#!/usr/bin/perl
use strict;
use warnings;
for (my $i = 0; $i < 64; $i++)
{
printf (" case %2d:\n", $i);
my $id4 = int ($i / 4);
my $im4 = int ($i % 4);
if ($im4 == 0)
{
printf (" w%d[%d] = 0;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 1)
{
printf (" w%d[%d] &= 0x000000ff;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 2)
{
printf (" w%d[%d] &= 0x0000ffff;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 3)
{
printf (" w%d[%d] &= 0x00ffffff;\n", $id4 / 4, $id4 % 4);
}
for (my $j = $id4 + 1; $j < 16; $j++)
{
my $jd4 = int ($j / 4);
my $jm4 = int ($j % 4);
printf (" w%d[%d] = 0;\n", $jd4, $jm4);
}
printf ("\n");
printf (" break;\n");
printf ("\n");
}

42
tools/code_generators/GEN_truncate_block_4x4_be.pl Normal file
View File

@ -0,0 +1,42 @@
#!/usr/bin/perl
use strict;
use warnings;
for (my $i = 0; $i < 16; $i++)
{
printf (" case %2d:\n", $i);
my $id4 = int ($i / 4);
my $im4 = int ($i % 4);
if ($im4 == 0)
{
printf (" w%d[%d] = 0;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 1)
{
printf (" w%d[%d] &= 0xff000000;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 2)
{
printf (" w%d[%d] &= 0xffff0000;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 3)
{
printf (" w%d[%d] &= 0xffffff00;\n", $id4 / 4, $id4 % 4);
}
for (my $j = $id4 + 1; $j < 4; $j++)
{
my $jd4 = int ($j / 4);
my $jm4 = int ($j % 4);
printf (" w%d[%d] = 0;\n", $jd4, $jm4);
}
printf ("\n");
printf (" break;\n");
printf ("\n");
}

42
tools/code_generators/GEN_truncate_block_4x4_le.pl Normal file
View File

@ -0,0 +1,42 @@
#!/usr/bin/perl
use strict;
use warnings;
for (my $i = 0; $i < 16; $i++)
{
printf (" case %2d:\n", $i);
my $id4 = int ($i / 4);
my $im4 = int ($i % 4);
if ($im4 == 0)
{
printf (" w%d[%d] = 0;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 1)
{
printf (" w%d[%d] &= 0x000000ff;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 2)
{
printf (" w%d[%d] &= 0x0000ffff;\n", $id4 / 4, $id4 % 4);
}
elsif ($im4 == 3)
{
printf (" w%d[%d] &= 0x00ffffff;\n", $id4 / 4, $id4 % 4);
}
for (my $j = $id4 + 1; $j < 4; $j++)
{
my $jd4 = int ($j / 4);
my $jm4 = int ($j % 4);
printf (" w%d[%d] = 0;\n", $jd4, $jm4);
}
printf ("\n");
printf (" break;\n");
printf ("\n");
}