arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-12-23 07:08:19 +00:00
Code Issues Releases Wiki Activity

Mix in pure kernel functions in various optimized kernels

Browse Source
This commit is contained in:
jsteube 2017-08-10 13:56:53 +02:00
parent 66f7590883
commit 68a8f70edb
10 changed files with 67 additions and 1234 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

149
OpenCL/m08400_a0-optimized.cl
View File

@ -13,6 +13,7 @@
#include "inc_rp.h"
#include "inc_rp.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#if VECT_SIZE == 1
#define uint_to_hex_lower8_le(i) (u32x) (l_bin2asc[(i)])
@ -26,134 +27,6 @@
#define uint_to_hex_lower8_le(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7], l_bin2asc[(i).s8], l_bin2asc[(i).s9], l_bin2asc[(i).sa], l_bin2asc[(i).sb], l_bin2asc[(i).sc], l_bin2asc[(i).sd], l_bin2asc[(i).se], l_bin2asc[(i).sf])
#endif
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{
u32x A = digest[0];
u32x B = digest[1];
u32x C = digest[2];
u32x D = digest[3];
u32x E = digest[4];
u32x w0_t = w0[0];
u32x w1_t = w0[1];
u32x w2_t = w0[2];
u32x w3_t = w0[3];
u32x w4_t = w1[0];
u32x w5_t = w1[1];
u32x w6_t = w1[2];
u32x w7_t = w1[3];
u32x w8_t = w2[0];
u32x w9_t = w2[1];
u32x wa_t = w2[2];
u32x wb_t = w2[3];
u32x wc_t = w3[0];
u32x wd_t = w3[1];
u32x we_t = w3[2];
u32x wf_t = w3[3];
#undef K
#define K SHA1C00
SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
#undef K
#define K SHA1C01
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
#undef K
#define K SHA1C02
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
#undef K
#define K SHA1C03
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
digest[0] += A;
digest[1] += B;
digest[2] += C;
digest[3] += D;
digest[4] += E;
}
__kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *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 void *esalt_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 u32 gid_max)
{
/**
@ -268,7 +141,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
u32x a;
u32x b;
@ -311,7 +184,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
w0[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -334,7 +207,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = 0;
w3[3] = 80 * 8;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
a = digest[0];
b = digest[1];
@ -371,7 +244,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
w0[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -394,7 +267,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = 0;
w3[3] = 80 * 8;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
COMPARE_M_SIMD (digest[3], digest[4], digest[2], digest[1]);
}
@ -534,7 +407,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
u32x a;
u32x b;
@ -577,7 +450,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
w0[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -600,7 +473,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = 0;
w3[3] = 80 * 8;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
a = digest[0];
b = digest[1];
@ -637,7 +510,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
w0[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -660,7 +533,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = 0;
w3[3] = 80 * 8;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
COMPARE_S_SIMD (digest[3], digest[4], digest[2], digest[1]);
}

149
OpenCL/m08400_a1-optimized.cl
View File

@ -11,6 +11,7 @@
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#if VECT_SIZE == 1
#define uint_to_hex_lower8_le(i) (u32x) (l_bin2asc[(i)])
@ -24,134 +25,6 @@
#define uint_to_hex_lower8_le(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7], l_bin2asc[(i).s8], l_bin2asc[(i).s9], l_bin2asc[(i).sa], l_bin2asc[(i).sb], l_bin2asc[(i).sc], l_bin2asc[(i).sd], l_bin2asc[(i).se], l_bin2asc[(i).sf])
#endif
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{
u32x A = digest[0];
u32x B = digest[1];
u32x C = digest[2];
u32x D = digest[3];
u32x E = digest[4];
u32x w0_t = w0[0];
u32x w1_t = w0[1];
u32x w2_t = w0[2];
u32x w3_t = w0[3];
u32x w4_t = w1[0];
u32x w5_t = w1[1];
u32x w6_t = w1[2];
u32x w7_t = w1[3];
u32x w8_t = w2[0];
u32x w9_t = w2[1];
u32x wa_t = w2[2];
u32x wb_t = w2[3];
u32x wc_t = w3[0];
u32x wd_t = w3[1];
u32x we_t = w3[2];
u32x wf_t = w3[3];
#undef K
#define K SHA1C00
SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
#undef K
#define K SHA1C01
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
#undef K
#define K SHA1C02
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
#undef K
#define K SHA1C03
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
digest[0] += A;
digest[1] += B;
digest[2] += C;
digest[3] += D;
digest[4] += E;
}
__kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *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 void *esalt_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 u32 gid_max)
{
/**
@ -324,7 +197,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
u32x a;
u32x b;
@ -367,7 +240,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
w0[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -390,7 +263,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = 0;
w3[3] = 80 * 8;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
a = digest[0];
b = digest[1];
@ -427,7 +300,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
w0[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -450,7 +323,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = 0;
w3[3] = 80 * 8;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
COMPARE_M_SIMD (digest[3], digest[4], digest[2], digest[1]);
}
@ -648,7 +521,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
u32x a;
u32x b;
@ -691,7 +564,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
w0[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -714,7 +587,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = 0;
w3[3] = 80 * 8;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
a = digest[0];
b = digest[1];
@ -751,7 +624,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
w0[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -774,7 +647,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w3[2] = 0;
w3[3] = 80 * 8;
sha1_transform (w0, w1, w2, w3, digest);
sha1_transform_vector (w0, w1, w2, w3, digest);
COMPARE_S_SIMD (digest[3], digest[4], digest[2], digest[1]);
}

149
OpenCL/m08400_a3-optimized.cl
View File

@ -11,6 +11,7 @@
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha1.cl"
#if VECT_SIZE == 1
#define uint_to_hex_lower8_le(i) (u32x) (l_bin2asc[(i)])
@ -24,134 +25,6 @@
#define uint_to_hex_lower8_le(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7], l_bin2asc[(i).s8], l_bin2asc[(i).s9], l_bin2asc[(i).sa], l_bin2asc[(i).sb], l_bin2asc[(i).sc], l_bin2asc[(i).sd], l_bin2asc[(i).se], l_bin2asc[(i).sf])
#endif
void sha1_transform (const u32x w0[4], const u32x w1[4], const u32x w2[4], const u32x w3[4], u32x digest[5])
{
u32x A = digest[0];
u32x B = digest[1];
u32x C = digest[2];
u32x D = digest[3];
u32x E = digest[4];
u32x w0_t = w0[0];
u32x w1_t = w0[1];
u32x w2_t = w0[2];
u32x w3_t = w0[3];
u32x w4_t = w1[0];
u32x w5_t = w1[1];
u32x w6_t = w1[2];
u32x w7_t = w1[3];
u32x w8_t = w2[0];
u32x w9_t = w2[1];
u32x wa_t = w2[2];
u32x wb_t = w2[3];
u32x wc_t = w3[0];
u32x wd_t = w3[1];
u32x we_t = w3[2];
u32x wf_t = w3[3];
#undef K
#define K SHA1C00
SHA1_STEP (SHA1_F0o, A, B, C, D, E, w0_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, w1_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, w2_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, w3_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, w4_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, w5_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, w6_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, w7_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, w8_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, w9_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, wa_t);
SHA1_STEP (SHA1_F0o, E, A, B, C, D, wb_t);
SHA1_STEP (SHA1_F0o, D, E, A, B, C, wc_t);
SHA1_STEP (SHA1_F0o, C, D, E, A, B, wd_t);
SHA1_STEP (SHA1_F0o, B, C, D, E, A, we_t);
SHA1_STEP (SHA1_F0o, A, B, C, D, E, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F0o, E, A, B, C, D, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F0o, D, E, A, B, C, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F0o, C, D, E, A, B, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F0o, B, C, D, E, A, w3_t);
#undef K
#define K SHA1C01
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w7_t);
#undef K
#define K SHA1C02
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F2o, A, B, C, D, E, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F2o, E, A, B, C, D, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F2o, D, E, A, B, C, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F2o, C, D, E, A, B, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F2o, B, C, D, E, A, wb_t);
#undef K
#define K SHA1C03
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, wf_t);
w0_t = rotl32 ((wd_t ^ w8_t ^ w2_t ^ w0_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w0_t);
w1_t = rotl32 ((we_t ^ w9_t ^ w3_t ^ w1_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w1_t);
w2_t = rotl32 ((wf_t ^ wa_t ^ w4_t ^ w2_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w2_t);
w3_t = rotl32 ((w0_t ^ wb_t ^ w5_t ^ w3_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w3_t);
w4_t = rotl32 ((w1_t ^ wc_t ^ w6_t ^ w4_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w4_t);
w5_t = rotl32 ((w2_t ^ wd_t ^ w7_t ^ w5_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, w5_t);
w6_t = rotl32 ((w3_t ^ we_t ^ w8_t ^ w6_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, w6_t);
w7_t = rotl32 ((w4_t ^ wf_t ^ w9_t ^ w7_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, w7_t);
w8_t = rotl32 ((w5_t ^ w0_t ^ wa_t ^ w8_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, w8_t);
w9_t = rotl32 ((w6_t ^ w1_t ^ wb_t ^ w9_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, w9_t);
wa_t = rotl32 ((w7_t ^ w2_t ^ wc_t ^ wa_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wa_t);
wb_t = rotl32 ((w8_t ^ w3_t ^ wd_t ^ wb_t), 1u); SHA1_STEP (SHA1_F1, A, B, C, D, E, wb_t);
wc_t = rotl32 ((w9_t ^ w4_t ^ we_t ^ wc_t), 1u); SHA1_STEP (SHA1_F1, E, A, B, C, D, wc_t);
wd_t = rotl32 ((wa_t ^ w5_t ^ wf_t ^ wd_t), 1u); SHA1_STEP (SHA1_F1, D, E, A, B, C, wd_t);
we_t = rotl32 ((wb_t ^ w6_t ^ w0_t ^ we_t), 1u); SHA1_STEP (SHA1_F1, C, D, E, A, B, we_t);
wf_t = rotl32 ((wc_t ^ w7_t ^ w1_t ^ wf_t), 1u); SHA1_STEP (SHA1_F1, B, C, D, E, A, wf_t);
digest[0] += A;
digest[1] += B;
digest[2] += C;
digest[3] += D;
digest[4] += E;
}
void m08400m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __global const bf_t *bfs_buf, __global void *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 void *esalt_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, __local u32 *l_bin2asc)
{
/**
@ -230,7 +103,7 @@ void m08400m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
u32x a;
u32x b;
@ -273,7 +146,7 @@ void m08400m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
w0_t[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -296,7 +169,7 @@ void m08400m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
w3_t[2] = 0;
w3_t[3] = (salt_len + 40) * 8;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
a = digest[0];
b = digest[1];
@ -333,7 +206,7 @@ void m08400m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
w0_t[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -356,7 +229,7 @@ void m08400m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
w3_t[2] = 0;
w3_t[3] = (salt_len + 40) * 8;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
COMPARE_M_SIMD (digest[3], digest[4], digest[2], digest[1]);
}
@ -452,7 +325,7 @@ void m08400s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
u32x a;
u32x b;
@ -495,7 +368,7 @@ void m08400s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
w0_t[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -518,7 +391,7 @@ void m08400s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
w3_t[2] = 0;
w3_t[3] = (salt_len + 40) * 8;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
a = digest[0];
b = digest[1];
@ -555,7 +428,7 @@ void m08400s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
digest[3] = SHA1M_D;
digest[4] = SHA1M_E;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
w0_t[0] = uint_to_hex_lower8_le ((d >> 16) & 255) << 0
| uint_to_hex_lower8_le ((d >> 24) & 255) << 16;
@ -578,7 +451,7 @@ void m08400s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
w3_t[2] = 0;
w3_t[3] = (salt_len + 40) * 8;
sha1_transform (w0_t, w1_t, w2_t, w3_t, digest);
sha1_transform_vector (w0_t, w1_t, w2_t, w3_t, digest);
COMPARE_S_SIMD (digest[3], digest[4], digest[2], digest[1]);
}

164
OpenCL/m13100_a0-optimized.cl
View File

@ -14,6 +14,8 @@
#include "inc_rp.h"
#include "inc_rp.cl"
#include "inc_simd.cl"
#include "inc_hash_md4.cl"
#include "inc_hash_md5.cl"
typedef struct
{
@ -138,168 +140,6 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[
return j;
}
void md4_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{
u32 a = digest[0];
u32 b = digest[1];
u32 c = digest[2];
u32 d = digest[3];
MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
}
void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{
u32 a = digest[0];
u32 b = digest[1];
u32 c = digest[2];
u32 d = digest[3];
u32 w0_t = w0[0];
u32 w1_t = w0[1];
u32 w2_t = w0[2];
u32 w3_t = w0[3];
u32 w4_t = w1[0];
u32 w5_t = w1[1];
u32 w6_t = w1[2];
u32 w7_t = w1[3];
u32 w8_t = w2[0];
u32 w9_t = w2[1];
u32 wa_t = w2[2];
u32 wb_t = w2[3];
u32 wc_t = w3[0];
u32 wd_t = w3[1];
u32 we_t = w3[2];
u32 wf_t = w3[3];
MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
}
void hmac_md5_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4])
{
w0[0] = w0[0] ^ 0x36363636;

164
OpenCL/m13100_a1-optimized.cl
View File

@ -12,6 +12,8 @@
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md4.cl"
#include "inc_hash_md5.cl"
typedef struct
{
@ -136,168 +138,6 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[
return j;
}
void md4_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{
u32 a = digest[0];
u32 b = digest[1];
u32 c = digest[2];
u32 d = digest[3];
MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
}
void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{
u32 a = digest[0];
u32 b = digest[1];
u32 c = digest[2];
u32 d = digest[3];
u32 w0_t = w0[0];
u32 w1_t = w0[1];
u32 w2_t = w0[2];
u32 w3_t = w0[3];
u32 w4_t = w1[0];
u32 w5_t = w1[1];
u32 w6_t = w1[2];
u32 w7_t = w1[3];
u32 w8_t = w2[0];
u32 w9_t = w2[1];
u32 wa_t = w2[2];
u32 wb_t = w2[3];
u32 wc_t = w3[0];
u32 wd_t = w3[1];
u32 we_t = w3[2];
u32 wf_t = w3[3];
MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
}
void hmac_md5_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4])
{
w0[0] = w0[0] ^ 0x36363636;

164
OpenCL/m13100_a3-optimized.cl
View File

@ -12,6 +12,8 @@
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_md4.cl"
#include "inc_hash_md5.cl"
typedef struct
{
@ -136,168 +138,6 @@ u8 rc4_next_16 (__local RC4_KEY *rc4_key, u8 i, u8 j, __global u32 *in, u32 out[
return j;
}
void md4_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{
u32 a = digest[0];
u32 b = digest[1];
u32 c = digest[2];
u32 d = digest[3];
MD4_STEP (MD4_Fo, a, b, c, d, w0[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w0[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w0[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w0[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w1[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w1[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w1[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w1[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w2[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w2[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w2[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w2[3], MD4C00, MD4S03);
MD4_STEP (MD4_Fo, a, b, c, d, w3[0], MD4C00, MD4S00);
MD4_STEP (MD4_Fo, d, a, b, c, w3[1], MD4C00, MD4S01);
MD4_STEP (MD4_Fo, c, d, a, b, w3[2], MD4C00, MD4S02);
MD4_STEP (MD4_Fo, b, c, d, a, w3[3], MD4C00, MD4S03);
MD4_STEP (MD4_Go, a, b, c, d, w0[0], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[0], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[0], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[0], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[1], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[1], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[1], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[1], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[2], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[2], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[2], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[2], MD4C01, MD4S13);
MD4_STEP (MD4_Go, a, b, c, d, w0[3], MD4C01, MD4S10);
MD4_STEP (MD4_Go, d, a, b, c, w1[3], MD4C01, MD4S11);
MD4_STEP (MD4_Go, c, d, a, b, w2[3], MD4C01, MD4S12);
MD4_STEP (MD4_Go, b, c, d, a, w3[3], MD4C01, MD4S13);
MD4_STEP (MD4_H , a, b, c, d, w0[0], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[0], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[0], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[0], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[2], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[2], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[2], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[2], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[1], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[1], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[1], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[1], MD4C02, MD4S23);
MD4_STEP (MD4_H , a, b, c, d, w0[3], MD4C02, MD4S20);
MD4_STEP (MD4_H , d, a, b, c, w2[3], MD4C02, MD4S21);
MD4_STEP (MD4_H , c, d, a, b, w1[3], MD4C02, MD4S22);
MD4_STEP (MD4_H , b, c, d, a, w3[3], MD4C02, MD4S23);
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
}
void md5_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[4])
{
u32 a = digest[0];
u32 b = digest[1];
u32 c = digest[2];
u32 d = digest[3];
u32 w0_t = w0[0];
u32 w1_t = w0[1];
u32 w2_t = w0[2];
u32 w3_t = w0[3];
u32 w4_t = w1[0];
u32 w5_t = w1[1];
u32 w6_t = w1[2];
u32 w7_t = w1[3];
u32 w8_t = w2[0];
u32 w9_t = w2[1];
u32 wa_t = w2[2];
u32 wb_t = w2[3];
u32 wc_t = w3[0];
u32 wd_t = w3[1];
u32 we_t = w3[2];
u32 wf_t = w3[3];
MD5_STEP (MD5_Fo, a, b, c, d, w0_t, MD5C00, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w1_t, MD5C01, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w2_t, MD5C02, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w3_t, MD5C03, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w4_t, MD5C04, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w5_t, MD5C05, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, w6_t, MD5C06, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, w7_t, MD5C07, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, w8_t, MD5C08, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, w9_t, MD5C09, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, wa_t, MD5C0a, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, wb_t, MD5C0b, MD5S03);
MD5_STEP (MD5_Fo, a, b, c, d, wc_t, MD5C0c, MD5S00);
MD5_STEP (MD5_Fo, d, a, b, c, wd_t, MD5C0d, MD5S01);
MD5_STEP (MD5_Fo, c, d, a, b, we_t, MD5C0e, MD5S02);
MD5_STEP (MD5_Fo, b, c, d, a, wf_t, MD5C0f, MD5S03);
MD5_STEP (MD5_Go, a, b, c, d, w1_t, MD5C10, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w6_t, MD5C11, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, wb_t, MD5C12, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w0_t, MD5C13, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w5_t, MD5C14, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, wa_t, MD5C15, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, wf_t, MD5C16, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w4_t, MD5C17, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, w9_t, MD5C18, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, we_t, MD5C19, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w3_t, MD5C1a, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, w8_t, MD5C1b, MD5S13);
MD5_STEP (MD5_Go, a, b, c, d, wd_t, MD5C1c, MD5S10);
MD5_STEP (MD5_Go, d, a, b, c, w2_t, MD5C1d, MD5S11);
MD5_STEP (MD5_Go, c, d, a, b, w7_t, MD5C1e, MD5S12);
MD5_STEP (MD5_Go, b, c, d, a, wc_t, MD5C1f, MD5S13);
MD5_STEP (MD5_H , a, b, c, d, w5_t, MD5C20, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w8_t, MD5C21, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, wb_t, MD5C22, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, we_t, MD5C23, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, w1_t, MD5C24, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w4_t, MD5C25, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, w7_t, MD5C26, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, wa_t, MD5C27, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, wd_t, MD5C28, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, w0_t, MD5C29, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, w3_t, MD5C2a, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, w6_t, MD5C2b, MD5S23);
MD5_STEP (MD5_H , a, b, c, d, w9_t, MD5C2c, MD5S20);
MD5_STEP (MD5_H , d, a, b, c, wc_t, MD5C2d, MD5S21);
MD5_STEP (MD5_H , c, d, a, b, wf_t, MD5C2e, MD5S22);
MD5_STEP (MD5_H , b, c, d, a, w2_t, MD5C2f, MD5S23);
MD5_STEP (MD5_I , a, b, c, d, w0_t, MD5C30, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w7_t, MD5C31, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, we_t, MD5C32, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w5_t, MD5C33, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, wc_t, MD5C34, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, w3_t, MD5C35, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, wa_t, MD5C36, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w1_t, MD5C37, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w8_t, MD5C38, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, wf_t, MD5C39, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w6_t, MD5C3a, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, wd_t, MD5C3b, MD5S33);
MD5_STEP (MD5_I , a, b, c, d, w4_t, MD5C3c, MD5S30);
MD5_STEP (MD5_I , d, a, b, c, wb_t, MD5C3d, MD5S31);
MD5_STEP (MD5_I , c, d, a, b, w2_t, MD5C3e, MD5S32);
MD5_STEP (MD5_I , b, c, d, a, w9_t, MD5C3f, MD5S33);
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
}
void hmac_md5_pad (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 ipad[4], u32 opad[4])
{
w0[0] = w0[0] ^ 0x36363636;

120
OpenCL/m13800_a0-optimized.cl
View File

@ -14,113 +14,11 @@
#include "inc_rp.h"
#include "inc_rp.cl"
#include "inc_simd.cl"
#include "inc_hash_sha256.cl"
__constant u32a k_sha256[64] =
void sha256_transform_transport_vector (const u32x w[16], u32x digest[8])
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,
SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
SHA256C10, SHA256C11, SHA256C12, SHA256C13,
SHA256C14, SHA256C15, SHA256C16, SHA256C17,
SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
SHA256C20, SHA256C21, SHA256C22, SHA256C23,
SHA256C24, SHA256C25, SHA256C26, SHA256C27,
SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
SHA256C30, SHA256C31, SHA256C32, SHA256C33,
SHA256C34, SHA256C35, SHA256C36, SHA256C37,
SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
};
void sha256_transform (const u32x w[16], u32x digest[8])
{
u32x a = digest[0];
u32x b = digest[1];
u32x c = digest[2];
u32x d = digest[3];
u32x e = digest[4];
u32x f = digest[5];
u32x g = digest[6];
u32x h = digest[7];
u32x w0_t = w[ 0];
u32x w1_t = w[ 1];
u32x w2_t = w[ 2];
u32x w3_t = w[ 3];
u32x w4_t = w[ 4];
u32x w5_t = w[ 5];
u32x w6_t = w[ 6];
u32x w7_t = w[ 7];
u32x w8_t = w[ 8];
u32x w9_t = w[ 9];
u32x wa_t = w[10];
u32x wb_t = w[11];
u32x wc_t = w[12];
u32x wd_t = w[13];
u32x we_t = w[14];
u32x wf_t = w[15];
#define ROUND_EXPAND() \
{ \
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); \
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); \
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); \
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); \
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); \
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); \
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); \
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); \
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); \
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); \
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); \
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); \
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); \
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); \
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); \
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); \
}
#define ROUND_STEP(i) \
{ \
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i + 0]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i + 1]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i + 2]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i + 3]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i + 4]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i + 5]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i + 6]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i + 7]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i + 8]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i + 9]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
}
ROUND_STEP (0);
#ifdef _unroll
#pragma unroll
#endif
for (int i = 16; i < 64; i += 16)
{
ROUND_EXPAND (); ROUND_STEP (i);
}
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
digest[4] += e;
digest[5] += f;
digest[6] += g;
digest[7] += h;
sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest);
}
void memcat64c_be (u32x block[16], const u32 offset, u32x carry[16])
@ -611,7 +509,7 @@ __kernel void m13800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[6] = SHA256M_G;
digest[7] = SHA256M_H;
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
w[ 0] = carry[ 0];
w[ 1] = carry[ 1];
@ -650,7 +548,7 @@ __kernel void m13800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
// we can always use pw_len here, since we add exactly the hash buffer size
memcat64c_be (w, out_len2, carry);
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
w[ 0] = carry[ 0];
w[ 1] = carry[ 1];
@ -675,7 +573,7 @@ __kernel void m13800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w[14] = 0;
w[15] = (out_len2 + 128) * 8;
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
const u32x d = digest[DGST_R0];
const u32x h = digest[DGST_R1];
@ -819,7 +717,7 @@ __kernel void m13800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[6] = SHA256M_G;
digest[7] = SHA256M_H;
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
w[ 0] = carry[ 0];
w[ 1] = carry[ 1];
@ -858,7 +756,7 @@ __kernel void m13800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
// we can always use pw_len here, since we add exactly the hash buffer size
memcat64c_be (w, out_len2, carry);
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
w[ 0] = carry[ 0];
w[ 1] = carry[ 1];
@ -883,7 +781,7 @@ __kernel void m13800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w[14] = 0;
w[15] = (out_len2 + 128) * 8;
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
const u32x d = digest[DGST_R0];
const u32x h = digest[DGST_R1];

120
OpenCL/m13800_a1-optimized.cl
View File

@ -12,113 +12,11 @@
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha256.cl"
__constant u32a k_sha256[64] =
void sha256_transform_transport_vector (const u32x w[16], u32x digest[8])
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,
SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
SHA256C10, SHA256C11, SHA256C12, SHA256C13,
SHA256C14, SHA256C15, SHA256C16, SHA256C17,
SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
SHA256C20, SHA256C21, SHA256C22, SHA256C23,
SHA256C24, SHA256C25, SHA256C26, SHA256C27,
SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
SHA256C30, SHA256C31, SHA256C32, SHA256C33,
SHA256C34, SHA256C35, SHA256C36, SHA256C37,
SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
};
void sha256_transform (const u32x w[16], u32x digest[8])
{
u32x a = digest[0];
u32x b = digest[1];
u32x c = digest[2];
u32x d = digest[3];
u32x e = digest[4];
u32x f = digest[5];
u32x g = digest[6];
u32x h = digest[7];
u32x w0_t = w[ 0];
u32x w1_t = w[ 1];
u32x w2_t = w[ 2];
u32x w3_t = w[ 3];
u32x w4_t = w[ 4];
u32x w5_t = w[ 5];
u32x w6_t = w[ 6];
u32x w7_t = w[ 7];
u32x w8_t = w[ 8];
u32x w9_t = w[ 9];
u32x wa_t = w[10];
u32x wb_t = w[11];
u32x wc_t = w[12];
u32x wd_t = w[13];
u32x we_t = w[14];
u32x wf_t = w[15];
#define ROUND_EXPAND() \
{ \
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); \
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); \
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); \
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); \
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); \
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); \
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); \
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); \
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); \
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); \
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); \
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); \
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); \
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); \
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); \
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); \
}
#define ROUND_STEP(i) \
{ \
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i + 0]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i + 1]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i + 2]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i + 3]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i + 4]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i + 5]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i + 6]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i + 7]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i + 8]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i + 9]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
}
ROUND_STEP (0);
#ifdef _unroll
#pragma unroll
#endif
for (int i = 16; i < 64; i += 16)
{
ROUND_EXPAND (); ROUND_STEP (i);
}
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
digest[4] += e;
digest[5] += f;
digest[6] += g;
digest[7] += h;
sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest);
}
void memcat64c_be (u32x block[16], const u32 offset, u32x carry[16])
@ -665,7 +563,7 @@ __kernel void m13800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[6] = SHA256M_G;
digest[7] = SHA256M_H;
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
w[ 0] = carry[ 0];
w[ 1] = carry[ 1];
@ -704,7 +602,7 @@ __kernel void m13800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
// we can always use pw_len here, since we add exactly the hash buffer size
memcat64c_be (w, pw_len2, carry);
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
w[ 0] = carry[ 0];
w[ 1] = carry[ 1];
@ -729,7 +627,7 @@ __kernel void m13800_m04 (__global pw_t *pws, __global const kernel_rule_t *rule
w[14] = 0;
w[15] = (pw_len2 + 128) * 8;
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
const u32x d = digest[DGST_R0];
const u32x h = digest[DGST_R1];
@ -929,7 +827,7 @@ __kernel void m13800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
digest[6] = SHA256M_G;
digest[7] = SHA256M_H;
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
w[ 0] = carry[ 0];
w[ 1] = carry[ 1];
@ -968,7 +866,7 @@ __kernel void m13800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
// we can always use pw_len here, since we add exactly the hash buffer size
memcat64c_be (w, pw_len2, carry);
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
w[ 0] = carry[ 0];
w[ 1] = carry[ 1];
@ -993,7 +891,7 @@ __kernel void m13800_s04 (__global pw_t *pws, __global const kernel_rule_t *rule
w[14] = 0;
w[15] = (pw_len2 + 128) * 8;
sha256_transform (w, digest);
sha256_transform_transport_vector (w, digest);
const u32x d = digest[DGST_R0];
const u32x h = digest[DGST_R1];

120
OpenCL/m13800_a3-optimized.cl
View File

@ -11,113 +11,11 @@
#include "inc_types.cl"
#include "inc_common.cl"
#include "inc_simd.cl"
#include "inc_hash_sha256.cl"
__constant u32a k_sha256[64] =
void sha256_transform_transport_vector (const u32x w[16], u32x digest[8])
{
SHA256C00, SHA256C01, SHA256C02, SHA256C03,
SHA256C04, SHA256C05, SHA256C06, SHA256C07,
SHA256C08, SHA256C09, SHA256C0a, SHA256C0b,
SHA256C0c, SHA256C0d, SHA256C0e, SHA256C0f,
SHA256C10, SHA256C11, SHA256C12, SHA256C13,
SHA256C14, SHA256C15, SHA256C16, SHA256C17,
SHA256C18, SHA256C19, SHA256C1a, SHA256C1b,
SHA256C1c, SHA256C1d, SHA256C1e, SHA256C1f,
SHA256C20, SHA256C21, SHA256C22, SHA256C23,
SHA256C24, SHA256C25, SHA256C26, SHA256C27,
SHA256C28, SHA256C29, SHA256C2a, SHA256C2b,
SHA256C2c, SHA256C2d, SHA256C2e, SHA256C2f,
SHA256C30, SHA256C31, SHA256C32, SHA256C33,
SHA256C34, SHA256C35, SHA256C36, SHA256C37,
SHA256C38, SHA256C39, SHA256C3a, SHA256C3b,
SHA256C3c, SHA256C3d, SHA256C3e, SHA256C3f,
};
void sha256_transform (const u32x w[16], u32x digest[8])
{
u32x a = digest[0];
u32x b = digest[1];
u32x c = digest[2];
u32x d = digest[3];
u32x e = digest[4];
u32x f = digest[5];
u32x g = digest[6];
u32x h = digest[7];
u32x w0_t = w[ 0];
u32x w1_t = w[ 1];
u32x w2_t = w[ 2];
u32x w3_t = w[ 3];
u32x w4_t = w[ 4];
u32x w5_t = w[ 5];
u32x w6_t = w[ 6];
u32x w7_t = w[ 7];
u32x w8_t = w[ 8];
u32x w9_t = w[ 9];
u32x wa_t = w[10];
u32x wb_t = w[11];
u32x wc_t = w[12];
u32x wd_t = w[13];
u32x we_t = w[14];
u32x wf_t = w[15];
#define ROUND_EXPAND() \
{ \
w0_t = SHA256_EXPAND (we_t, w9_t, w1_t, w0_t); \
w1_t = SHA256_EXPAND (wf_t, wa_t, w2_t, w1_t); \
w2_t = SHA256_EXPAND (w0_t, wb_t, w3_t, w2_t); \
w3_t = SHA256_EXPAND (w1_t, wc_t, w4_t, w3_t); \
w4_t = SHA256_EXPAND (w2_t, wd_t, w5_t, w4_t); \
w5_t = SHA256_EXPAND (w3_t, we_t, w6_t, w5_t); \
w6_t = SHA256_EXPAND (w4_t, wf_t, w7_t, w6_t); \
w7_t = SHA256_EXPAND (w5_t, w0_t, w8_t, w7_t); \
w8_t = SHA256_EXPAND (w6_t, w1_t, w9_t, w8_t); \
w9_t = SHA256_EXPAND (w7_t, w2_t, wa_t, w9_t); \
wa_t = SHA256_EXPAND (w8_t, w3_t, wb_t, wa_t); \
wb_t = SHA256_EXPAND (w9_t, w4_t, wc_t, wb_t); \
wc_t = SHA256_EXPAND (wa_t, w5_t, wd_t, wc_t); \
wd_t = SHA256_EXPAND (wb_t, w6_t, we_t, wd_t); \
we_t = SHA256_EXPAND (wc_t, w7_t, wf_t, we_t); \
wf_t = SHA256_EXPAND (wd_t, w8_t, w0_t, wf_t); \
}
#define ROUND_STEP(i) \
{ \
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w0_t, k_sha256[i + 0]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w1_t, k_sha256[i + 1]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, w2_t, k_sha256[i + 2]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, w3_t, k_sha256[i + 3]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, w4_t, k_sha256[i + 4]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, w5_t, k_sha256[i + 5]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, w6_t, k_sha256[i + 6]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, w7_t, k_sha256[i + 7]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, a, b, c, d, e, f, g, h, w8_t, k_sha256[i + 8]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, h, a, b, c, d, e, f, g, w9_t, k_sha256[i + 9]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, g, h, a, b, c, d, e, f, wa_t, k_sha256[i + 10]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, f, g, h, a, b, c, d, e, wb_t, k_sha256[i + 11]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, e, f, g, h, a, b, c, d, wc_t, k_sha256[i + 12]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, d, e, f, g, h, a, b, c, wd_t, k_sha256[i + 13]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, c, d, e, f, g, h, a, b, we_t, k_sha256[i + 14]); \
SHA256_STEP (SHA256_F0o, SHA256_F1o, b, c, d, e, f, g, h, a, wf_t, k_sha256[i + 15]); \
}
ROUND_STEP (0);
#ifdef _unroll
#pragma unroll
#endif
for (int i = 16; i < 64; i += 16)
{
ROUND_EXPAND (); ROUND_STEP (i);
}
digest[0] += a;
digest[1] += b;
digest[2] += c;
digest[3] += d;
digest[4] += e;
digest[5] += f;
digest[6] += g;
digest[7] += h;
sha256_transform_vector (w + 0, w + 4, w + 8, w + 12, digest);
}
void memcat64c_be (u32x block[16], const u32 offset, u32x carry[16])
@ -568,7 +466,7 @@ void m13800m (__local u32 *s_esalt, u32 w[16], const u32 pw_len, __global pw_t *
digest[6] = SHA256M_G;
digest[7] = SHA256M_H;
sha256_transform (w_t, digest);
sha256_transform_transport_vector (w_t, digest);
w_t[ 0] = carry[ 0];
w_t[ 1] = carry[ 1];
@ -607,7 +505,7 @@ void m13800m (__local u32 *s_esalt, u32 w[16], const u32 pw_len, __global pw_t *
// we can always use pw_len here, since we add exactly the hash buffer size
memcat64c_be (w_t, pw_len, carry);
sha256_transform (w_t, digest);
sha256_transform_transport_vector (w_t, digest);
w_t[ 0] = carry[ 0];
w_t[ 1] = carry[ 1];
@ -632,7 +530,7 @@ void m13800m (__local u32 *s_esalt, u32 w[16], const u32 pw_len, __global pw_t *
w_t[14] = 0;
w_t[15] = (pw_len + 128) * 8;
sha256_transform (w_t, digest);
sha256_transform_transport_vector (w_t, digest);
const u32x d = digest[DGST_R0];
const u32x h = digest[DGST_R1];
@ -728,7 +626,7 @@ void m13800s (__local u32 *s_esalt, u32 w[16], const u32 pw_len, __global pw_t *
digest[6] = SHA256M_G;
digest[7] = SHA256M_H;
sha256_transform (w_t, digest);
sha256_transform_transport_vector (w_t, digest);
w_t[ 0] = carry[ 0];
w_t[ 1] = carry[ 1];
@ -767,7 +665,7 @@ void m13800s (__local u32 *s_esalt, u32 w[16], const u32 pw_len, __global pw_t *
// we can always use pw_len here, since we add exactly the hash buffer size
memcat64c_be (w_t, pw_len, carry);
sha256_transform (w_t, digest);
sha256_transform_transport_vector (w_t, digest);
w_t[ 0] = carry[ 0];
w_t[ 1] = carry[ 1];
@ -792,7 +690,7 @@ void m13800s (__local u32 *s_esalt, u32 w[16], const u32 pw_len, __global pw_t *
w_t[14] = 0;
w_t[15] = (pw_len + 128) * 8;
sha256_transform (w_t, digest);
sha256_transform_transport_vector (w_t, digest);
const u32x d = digest[DGST_R0];
const u32x h = digest[DGST_R1];

2
tools/test.sh
View File

@ -24,7 +24,7 @@ NEVER_CRACK="11600 14900"
SLOW_ALGOS="400 500 501 1600 1800 2100 2500 3200 5200 5800 6211 6212 6213 6221 6222 6223 6231 6232 6233 6241 6242 6243 6251 6261 6271 6281 6300 6400 6500 6600 6700 6800 7100 7200 7400 7900 8200 8800 8900 9000 9100 9200 9300 9400 9500 9600 10000 10300 10500 10700 10900 11300 11600 11900 12000 12001 12100 12200 12300 12400 12500 12700 12800 12900 13000 13200 13400 13600 14600 14700 14800 15100 15200 15300 15600 15700 15800"
OPTS="--quiet --force --potfile-disable --runtime 400 --gpu-temp-disable --weak-hash-threshold=0"
OPTS="--quiet --force --potfile-disable --runtime 400 --gpu-temp-disable --weak-hash-threshold=0 -O"
OUTD="test_$(date +%s)"