mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-26 09:58:16 +00:00
Fix vector function calls and datatypes
This commit is contained in:
jsteube
parent
1eb249c5b4
commit
2802f1d592
@ -82,13 +82,13 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
md5_hmac_ctx_vector_t ctx;
|
||||
md5_hmac_ctx_t ctx;
|
||||
|
||||
md5_hmac_init_vector (&ctx, c, pw_len + comb_len);
|
||||
md5_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
md5_hmac_update_vector (&ctx, s, salt_len);
|
||||
md5_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
md5_hmac_final_vector (&ctx);
|
||||
md5_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
@ -180,13 +180,13 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
md5_hmac_ctx_vector_t ctx;
|
||||
md5_hmac_ctx_t ctx;
|
||||
|
||||
md5_hmac_init_vector (&ctx, c, pw_len + comb_len);
|
||||
md5_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
md5_hmac_update_vector (&ctx, s, salt_len);
|
||||
md5_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
md5_hmac_final_vector (&ctx);
|
||||
md5_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
|
||||
@ -82,13 +82,13 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
sha1_hmac_ctx_vector_t ctx;
|
||||
sha1_hmac_ctx_t ctx;
|
||||
|
||||
sha1_hmac_init_vector (&ctx, c, pw_len + comb_len);
|
||||
sha1_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
sha1_hmac_update_vector (&ctx, s, salt_len);
|
||||
sha1_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
sha1_hmac_final_vector (&ctx);
|
||||
sha1_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
@ -180,13 +180,13 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
sha1_hmac_ctx_vector_t ctx;
|
||||
sha1_hmac_ctx_t ctx;
|
||||
|
||||
sha1_hmac_init_vector (&ctx, c, pw_len + comb_len);
|
||||
sha1_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
sha1_hmac_update_vector (&ctx, s, salt_len);
|
||||
sha1_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
sha1_hmac_final_vector (&ctx);
|
||||
sha1_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
|
||||
@ -47,7 +47,7 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -140,7 +140,7 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -32,7 +32,7 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -114,7 +114,7 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -82,13 +82,13 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
sha256_hmac_ctx_vector_t ctx;
|
||||
sha256_hmac_ctx_t ctx;
|
||||
|
||||
sha256_hmac_init_vector (&ctx, c, pw_len + comb_len);
|
||||
sha256_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
sha256_hmac_update_vector (&ctx, s, salt_len);
|
||||
sha256_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
sha256_hmac_final_vector (&ctx);
|
||||
sha256_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
@ -180,13 +180,13 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
sha256_hmac_ctx_vector_t ctx;
|
||||
sha256_hmac_ctx_t ctx;
|
||||
|
||||
sha256_hmac_init_vector (&ctx, c, pw_len + comb_len);
|
||||
sha256_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
sha256_hmac_update_vector (&ctx, s, salt_len);
|
||||
sha256_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
sha256_hmac_final_vector (&ctx);
|
||||
sha256_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
|
||||
@ -82,13 +82,13 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
sha512_hmac_ctx_vector_t ctx;
|
||||
sha512_hmac_ctx_t ctx;
|
||||
|
||||
sha512_hmac_init_vector (&ctx, c, pw_len + comb_len);
|
||||
sha512_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
sha512_hmac_update_vector (&ctx, s, salt_len);
|
||||
sha512_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
sha512_hmac_final_vector (&ctx);
|
||||
sha512_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = l32_from_64_S (ctx.opad.h[7]);
|
||||
const u32 r1 = h32_from_64_S (ctx.opad.h[7]);
|
||||
@ -180,13 +180,13 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
c[i] |= w[i];
|
||||
}
|
||||
|
||||
sha512_hmac_ctx_vector_t ctx;
|
||||
sha512_hmac_ctx_t ctx;
|
||||
|
||||
sha512_hmac_init_vector (&ctx, c, pw_len + comb_len);
|
||||
sha512_hmac_init (&ctx, c, pw_len + comb_len);
|
||||
|
||||
sha512_hmac_update_vector (&ctx, s, salt_len);
|
||||
sha512_hmac_update (&ctx, s, salt_len);
|
||||
|
||||
sha512_hmac_final_vector (&ctx);
|
||||
sha512_hmac_final (&ctx);
|
||||
|
||||
const u32 r0 = l32_from_64_S (ctx.opad.h[7]);
|
||||
const u32 r1 = h32_from_64_S (ctx.opad.h[7]);
|
||||
|
||||
@ -77,7 +77,7 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -205,7 +205,7 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -62,7 +62,7 @@ __kernel void m02610_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -179,7 +179,7 @@ __kernel void m02610_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -77,7 +77,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -228,7 +228,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -62,7 +62,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -202,7 +202,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -25,7 +25,7 @@
|
||||
#define uint_to_hex_lower8(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
|
||||
|
||||
__kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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)
|
||||
__kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32 *words_buf_r, __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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
@ -62,7 +62,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32x w[64] = { 0 };
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
@ -75,7 +75,7 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -88,13 +88,13 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* loop
|
||||
*/
|
||||
|
||||
u32x w0l = w[0];
|
||||
u32 w0l = w[0];
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
|
||||
const u32 w0r = words_buf_r[il_pos / VECT_SIZE];
|
||||
|
||||
const u32x w0lr = w0l | w0r;
|
||||
const u32 w0lr = w0l | w0r;
|
||||
|
||||
w[0] = w0lr;
|
||||
|
||||
@ -106,10 +106,10 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_final_vector (&ctx0);
|
||||
|
||||
const u32x a = ctx0.h[0];
|
||||
const u32x b = ctx0.h[1];
|
||||
const u32x c = ctx0.h[2];
|
||||
const u32x d = ctx0.h[3];
|
||||
const u32 a = ctx0.h[0];
|
||||
const u32 b = ctx0.h[1];
|
||||
const u32 c = ctx0.h[2];
|
||||
const u32 d = ctx0.h[3];
|
||||
|
||||
md5_ctx_vector_t ctx;
|
||||
|
||||
@ -161,16 +161,16 @@ __kernel void m02810_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_transform_vector (ctx.w0, ctx.w1, ctx.w2, ctx.w3, ctx.h);
|
||||
|
||||
const u32x r0 = ctx.h[DGST_R0];
|
||||
const u32x r1 = ctx.h[DGST_R1];
|
||||
const u32x r2 = ctx.h[DGST_R2];
|
||||
const u32x r3 = ctx.h[DGST_R3];
|
||||
const u32 r0 = ctx.h[DGST_R0];
|
||||
const u32 r1 = ctx.h[DGST_R1];
|
||||
const u32 r2 = ctx.h[DGST_R2];
|
||||
const u32 r3 = ctx.h[DGST_R3];
|
||||
|
||||
COMPARE_M_SIMD (r0, r1, r2, r3);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32x *words_buf_r, __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)
|
||||
__kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const pw_t *combs_buf, __constant const u32 *words_buf_r, __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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
@ -219,7 +219,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 pw_lenv = ceil ((float) pw_len / 4);
|
||||
|
||||
u32x w[64] = { 0 };
|
||||
u32 w[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
@ -232,7 +232,7 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -245,13 +245,13 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
* loop
|
||||
*/
|
||||
|
||||
u32x w0l = w[0];
|
||||
u32 w0l = w[0];
|
||||
|
||||
for (u32 il_pos = 0; il_pos < il_cnt; il_pos += VECT_SIZE)
|
||||
{
|
||||
const u32x w0r = words_buf_r[il_pos / VECT_SIZE];
|
||||
const u32 w0r = words_buf_r[il_pos / VECT_SIZE];
|
||||
|
||||
const u32x w0lr = w0l | w0r;
|
||||
const u32 w0lr = w0l | w0r;
|
||||
|
||||
w[0] = w0lr;
|
||||
|
||||
@ -263,10 +263,10 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_final_vector (&ctx0);
|
||||
|
||||
const u32x a = ctx0.h[0];
|
||||
const u32x b = ctx0.h[1];
|
||||
const u32x c = ctx0.h[2];
|
||||
const u32x d = ctx0.h[3];
|
||||
const u32 a = ctx0.h[0];
|
||||
const u32 b = ctx0.h[1];
|
||||
const u32 c = ctx0.h[2];
|
||||
const u32 d = ctx0.h[3];
|
||||
|
||||
md5_ctx_vector_t ctx;
|
||||
|
||||
@ -318,10 +318,10 @@ __kernel void m02810_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_transform_vector (ctx.w0, ctx.w1, ctx.w2, ctx.w3, ctx.h);
|
||||
|
||||
const u32x r0 = ctx.h[DGST_R0];
|
||||
const u32x r1 = ctx.h[DGST_R1];
|
||||
const u32x r2 = ctx.h[DGST_R2];
|
||||
const u32x r3 = ctx.h[DGST_R3];
|
||||
const u32 r0 = ctx.h[DGST_R0];
|
||||
const u32 r1 = ctx.h[DGST_R1];
|
||||
const u32 r2 = ctx.h[DGST_R2];
|
||||
const u32 r3 = ctx.h[DGST_R3];
|
||||
|
||||
COMPARE_S_SIMD (r0, r1, r2, r3);
|
||||
}
|
||||
|
||||
@ -77,7 +77,7 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -218,7 +218,7 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -62,7 +62,7 @@ __kernel void m03710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -192,7 +192,7 @@ __kernel void m03710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -77,7 +77,7 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -228,7 +228,7 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -62,7 +62,7 @@ __kernel void m03910_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -202,7 +202,7 @@ __kernel void m03910_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[8] = { 0 };
|
||||
u32 s[8] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -77,7 +77,7 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -205,7 +205,7 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -62,7 +62,7 @@ __kernel void m04310_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
@ -179,7 +179,7 @@ __kernel void m04310_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
const u32 salt_lenv = ceil ((float) salt_len / 4);
|
||||
|
||||
u32x s[64] = { 0 };
|
||||
u32 s[64] = { 0 };
|
||||
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
|
||||
@ -72,7 +72,7 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_update (&ctx, w, pw_len);
|
||||
|
||||
md5_update_vector (&ctx, s, salt_len);
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
@ -154,7 +154,7 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_update (&ctx, w, pw_len);
|
||||
|
||||
md5_update_vector (&ctx, s, salt_len);
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
|
||||
@ -57,7 +57,7 @@ __kernel void m04800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
|
||||
|
||||
md5_update_vector (&ctx, s, salt_len);
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
@ -126,7 +126,7 @@ __kernel void m04800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_update_global (&ctx, combs_buf[il_pos].i, combs_buf[il_pos].pw_len);
|
||||
|
||||
md5_update_vector (&ctx, s, salt_len);
|
||||
md5_update (&ctx, s, salt_len);
|
||||
|
||||
md5_final (&ctx);
|
||||
|
||||
|
||||
@ -89,10 +89,10 @@ __kernel void m05400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha1_hmac_final (&ctx);
|
||||
|
||||
const u32x r0 = ctx.opad.h[DGST_R0];
|
||||
const u32x r1 = ctx.opad.h[DGST_R1];
|
||||
const u32x r2 = ctx.opad.h[DGST_R2];
|
||||
const u32x r3 = ctx.opad.h[DGST_R3];
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
const u32 r2 = ctx.opad.h[DGST_R2];
|
||||
const u32 r3 = ctx.opad.h[DGST_R3];
|
||||
|
||||
COMPARE_M_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
@ -184,10 +184,10 @@ __kernel void m05400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha1_hmac_final (&ctx);
|
||||
|
||||
const u32x r0 = ctx.opad.h[DGST_R0];
|
||||
const u32x r1 = ctx.opad.h[DGST_R1];
|
||||
const u32x r2 = ctx.opad.h[DGST_R2];
|
||||
const u32x r3 = ctx.opad.h[DGST_R3];
|
||||
const u32 r0 = ctx.opad.h[DGST_R0];
|
||||
const u32 r1 = ctx.opad.h[DGST_R1];
|
||||
const u32 r2 = ctx.opad.h[DGST_R2];
|
||||
const u32 r3 = ctx.opad.h[DGST_R3];
|
||||
|
||||
COMPARE_S_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
|
||||
@ -344,18 +344,18 @@ __constant u32a c_skb[8][64] =
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
|
||||
#endif
|
||||
|
||||
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x r = data[0];
|
||||
u32x l = data[1];
|
||||
u32 r = data[0];
|
||||
u32 l = data[1];
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
u32 u;
|
||||
u32 t;
|
||||
|
||||
u = Kc[i + 0] ^ rotl32 (r, 30u);
|
||||
t = Kd[i + 0] ^ rotl32 (r, 26u);
|
||||
@ -386,9 +386,9 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
|
||||
iv[1] = r;
|
||||
}
|
||||
|
||||
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
|
||||
void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64])
|
||||
{
|
||||
u32x tt;
|
||||
u32 tt;
|
||||
|
||||
PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
|
||||
HPERM_OP (c, tt, 2, 0xcccc0000);
|
||||
@ -423,13 +423,13 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
c = c & 0x0fffffff;
|
||||
d = d & 0x0fffffff;
|
||||
|
||||
const u32x c00 = (c >> 0) & 0x0000003f;
|
||||
const u32x c06 = (c >> 6) & 0x00383003;
|
||||
const u32x c07 = (c >> 7) & 0x0000003c;
|
||||
const u32x c13 = (c >> 13) & 0x0000060f;
|
||||
const u32x c20 = (c >> 20) & 0x00000001;
|
||||
const u32 c00 = (c >> 0) & 0x0000003f;
|
||||
const u32 c06 = (c >> 6) & 0x00383003;
|
||||
const u32 c07 = (c >> 7) & 0x0000003c;
|
||||
const u32 c13 = (c >> 13) & 0x0000060f;
|
||||
const u32 c20 = (c >> 20) & 0x00000001;
|
||||
|
||||
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
| BOX (((c06 >> 0) & 0xff)
|
||||
|((c07 >> 0) & 0xff), 1, s_skb)
|
||||
| BOX (((c13 >> 0) & 0xff)
|
||||
@ -438,12 +438,12 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
|((c13 >> 8) & 0xff)
|
||||
|((c06 >> 16) & 0xff), 3, s_skb);
|
||||
|
||||
const u32x d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32x d07 = (d >> 7) & 0x00003f03;
|
||||
const u32x d21 = (d >> 21) & 0x0000000f;
|
||||
const u32x d22 = (d >> 22) & 0x00000030;
|
||||
const u32 d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32 d07 = (d >> 7) & 0x00003f03;
|
||||
const u32 d21 = (d >> 21) & 0x0000000f;
|
||||
const u32 d22 = (d >> 22) & 0x00000030;
|
||||
|
||||
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
| BOX (((d07 >> 0) & 0xff)
|
||||
|((d00 >> 8) & 0xff), 5, s_skb)
|
||||
| BOX (((d07 >> 8) & 0xff), 6, s_skb)
|
||||
@ -455,9 +455,9 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
}
|
||||
}
|
||||
|
||||
void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
|
||||
void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
|
||||
{
|
||||
u32x t[8];
|
||||
u32 t[8];
|
||||
|
||||
t[0] = (w0 >> 0) & 0xff;
|
||||
t[1] = (w0 >> 8) & 0xff;
|
||||
@ -468,7 +468,7 @@ void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
|
||||
t[6] = (w1 >> 16) & 0xff;
|
||||
t[7] = (w1 >> 24) & 0xff;
|
||||
|
||||
u32x k[8];
|
||||
u32 k[8];
|
||||
|
||||
k[0] = (t[0] >> 0);
|
||||
k[1] = (t[0] << 7) | (t[1] >> 1);
|
||||
|
||||
@ -341,18 +341,18 @@ __constant u32a c_skb[8][64] =
|
||||
#define BOX(i,n,S) (u32x) ((S)[(n)][(i).s0], (S)[(n)][(i).s1], (S)[(n)][(i).s2], (S)[(n)][(i).s3], (S)[(n)][(i).s4], (S)[(n)][(i).s5], (S)[(n)][(i).s6], (S)[(n)][(i).s7], (S)[(n)][(i).s8], (S)[(n)][(i).s9], (S)[(n)][(i).sa], (S)[(n)][(i).sb], (S)[(n)][(i).sc], (S)[(n)][(i).sd], (S)[(n)][(i).se], (S)[(n)][(i).sf])
|
||||
#endif
|
||||
|
||||
void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
void _des_crypt_encrypt (u32 iv[2], u32 data[2], u32 Kc[16], u32 Kd[16], __local u32 (*s_SPtrans)[64])
|
||||
{
|
||||
u32x r = data[0];
|
||||
u32x l = data[1];
|
||||
u32 r = data[0];
|
||||
u32 l = data[1];
|
||||
|
||||
#ifdef _unroll
|
||||
#pragma unroll
|
||||
#endif
|
||||
for (u32 i = 0; i < 16; i += 2)
|
||||
{
|
||||
u32x u;
|
||||
u32x t;
|
||||
u32 u;
|
||||
u32 t;
|
||||
|
||||
u = Kc[i + 0] ^ rotl32 (r, 30u);
|
||||
t = Kd[i + 0] ^ rotl32 (r, 26u);
|
||||
@ -383,9 +383,9 @@ void _des_crypt_encrypt (u32x iv[2], u32x data[2], u32x Kc[16], u32x Kd[16], __l
|
||||
iv[1] = r;
|
||||
}
|
||||
|
||||
void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32 (*s_skb)[64])
|
||||
void _des_crypt_keysetup (u32 c, u32 d, u32 Kc[16], u32 Kd[16], __local u32 (*s_skb)[64])
|
||||
{
|
||||
u32x tt;
|
||||
u32 tt;
|
||||
|
||||
PERM_OP (d, c, tt, 4, 0x0f0f0f0f);
|
||||
HPERM_OP (c, tt, 2, 0xcccc0000);
|
||||
@ -420,13 +420,13 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
c = c & 0x0fffffff;
|
||||
d = d & 0x0fffffff;
|
||||
|
||||
const u32x c00 = (c >> 0) & 0x0000003f;
|
||||
const u32x c06 = (c >> 6) & 0x00383003;
|
||||
const u32x c07 = (c >> 7) & 0x0000003c;
|
||||
const u32x c13 = (c >> 13) & 0x0000060f;
|
||||
const u32x c20 = (c >> 20) & 0x00000001;
|
||||
const u32 c00 = (c >> 0) & 0x0000003f;
|
||||
const u32 c06 = (c >> 6) & 0x00383003;
|
||||
const u32 c07 = (c >> 7) & 0x0000003c;
|
||||
const u32 c13 = (c >> 13) & 0x0000060f;
|
||||
const u32 c20 = (c >> 20) & 0x00000001;
|
||||
|
||||
u32x s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
u32 s = BOX (((c00 >> 0) & 0xff), 0, s_skb)
|
||||
| BOX (((c06 >> 0) & 0xff)
|
||||
|((c07 >> 0) & 0xff), 1, s_skb)
|
||||
| BOX (((c13 >> 0) & 0xff)
|
||||
@ -435,12 +435,12 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
|((c13 >> 8) & 0xff)
|
||||
|((c06 >> 16) & 0xff), 3, s_skb);
|
||||
|
||||
const u32x d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32x d07 = (d >> 7) & 0x00003f03;
|
||||
const u32x d21 = (d >> 21) & 0x0000000f;
|
||||
const u32x d22 = (d >> 22) & 0x00000030;
|
||||
const u32 d00 = (d >> 0) & 0x00003c3f;
|
||||
const u32 d07 = (d >> 7) & 0x00003f03;
|
||||
const u32 d21 = (d >> 21) & 0x0000000f;
|
||||
const u32 d22 = (d >> 22) & 0x00000030;
|
||||
|
||||
u32x t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
u32 t = BOX (((d00 >> 0) & 0xff), 4, s_skb)
|
||||
| BOX (((d07 >> 0) & 0xff)
|
||||
|((d00 >> 8) & 0xff), 5, s_skb)
|
||||
| BOX (((d07 >> 8) & 0xff), 6, s_skb)
|
||||
@ -452,9 +452,9 @@ void _des_crypt_keysetup (u32x c, u32x d, u32x Kc[16], u32x Kd[16], __local u32
|
||||
}
|
||||
}
|
||||
|
||||
void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
|
||||
void transform_netntlmv1_key (const u32 w0, const u32 w1, u32 out[2])
|
||||
{
|
||||
u32x t[8];
|
||||
u32 t[8];
|
||||
|
||||
t[0] = (w0 >> 0) & 0xff;
|
||||
t[1] = (w0 >> 8) & 0xff;
|
||||
@ -465,7 +465,7 @@ void transform_netntlmv1_key (const u32x w0, const u32x w1, u32x out[2])
|
||||
t[6] = (w1 >> 16) & 0xff;
|
||||
t[7] = (w1 >> 24) & 0xff;
|
||||
|
||||
u32x k[8];
|
||||
u32 k[8];
|
||||
|
||||
k[0] = (t[0] >> 0);
|
||||
k[1] = (t[0] << 7) | (t[1] >> 1);
|
||||
|
||||
@ -59,10 +59,10 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha384_final (&ctx);
|
||||
|
||||
const u32x r0 = l32_from_64 (ctx.h[3]);
|
||||
const u32x r1 = h32_from_64 (ctx.h[3]);
|
||||
const u32x r2 = l32_from_64 (ctx.h[2]);
|
||||
const u32x r3 = h32_from_64 (ctx.h[2]);
|
||||
const u32 r0 = l32_from_64 (ctx.h[3]);
|
||||
const u32 r1 = h32_from_64 (ctx.h[3]);
|
||||
const u32 r2 = l32_from_64 (ctx.h[2]);
|
||||
const u32 r3 = h32_from_64 (ctx.h[2]);
|
||||
|
||||
COMPARE_M_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
@ -124,10 +124,10 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha384_final (&ctx);
|
||||
|
||||
const u32x r0 = l32_from_64 (ctx.h[3]);
|
||||
const u32x r1 = h32_from_64 (ctx.h[3]);
|
||||
const u32x r2 = l32_from_64 (ctx.h[2]);
|
||||
const u32x r3 = h32_from_64 (ctx.h[2]);
|
||||
const u32 r0 = l32_from_64 (ctx.h[3]);
|
||||
const u32 r1 = h32_from_64 (ctx.h[3]);
|
||||
const u32 r2 = l32_from_64 (ctx.h[2]);
|
||||
const u32 r3 = h32_from_64 (ctx.h[2]);
|
||||
|
||||
COMPARE_S_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
|
||||
@ -46,10 +46,10 @@ __kernel void m10800_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha384_final (&ctx);
|
||||
|
||||
const u32x r0 = l32_from_64 (ctx.h[3]);
|
||||
const u32x r1 = h32_from_64 (ctx.h[3]);
|
||||
const u32x r2 = l32_from_64 (ctx.h[2]);
|
||||
const u32x r3 = h32_from_64 (ctx.h[2]);
|
||||
const u32 r0 = l32_from_64 (ctx.h[3]);
|
||||
const u32 r1 = h32_from_64 (ctx.h[3]);
|
||||
const u32 r2 = l32_from_64 (ctx.h[2]);
|
||||
const u32 r3 = h32_from_64 (ctx.h[2]);
|
||||
|
||||
COMPARE_M_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
@ -100,10 +100,10 @@ __kernel void m10800_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha384_final (&ctx);
|
||||
|
||||
const u32x r0 = l32_from_64 (ctx.h[3]);
|
||||
const u32x r1 = h32_from_64 (ctx.h[3]);
|
||||
const u32x r2 = l32_from_64 (ctx.h[2]);
|
||||
const u32x r3 = h32_from_64 (ctx.h[2]);
|
||||
const u32 r0 = l32_from_64 (ctx.h[3]);
|
||||
const u32 r1 = h32_from_64 (ctx.h[3]);
|
||||
const u32 r2 = l32_from_64 (ctx.h[2]);
|
||||
const u32 r3 = h32_from_64 (ctx.h[2]);
|
||||
|
||||
COMPARE_S_SCALAR (r0, r1, r2, r3);
|
||||
}
|
||||
|
||||
Loading…
Reference in New Issue
Block a user