mirror of
https://github.com/hashcat/hashcat.git
synced 2024-11-22 16:18:09 +00:00
Fix some datatypes
This commit is contained in:
parent
5c75eb84b0
commit
0334cdb277
@ -2879,7 +2879,7 @@ void sha384_hmac_update_vector (sha384_hmac_ctx_vector_t *ctx, const u32x *w, co
|
||||
sha384_update_vector (&ctx->ipad, w, len);
|
||||
}
|
||||
|
||||
void sha384_hmac_final_vector (sha384_hmac_ctx_t *ctx)
|
||||
void sha384_hmac_final_vector (sha384_hmac_ctx_vector_t *ctx)
|
||||
{
|
||||
sha384_final_vector (&ctx->ipad);
|
||||
|
||||
|
@ -2396,10 +2396,10 @@ void sha512_update_vector_utf16le (sha512_ctx_vector_t *ctx, const u32x *w, cons
|
||||
w3[2] = w[pos4 + 14];
|
||||
w3[3] = w[pos4 + 15];
|
||||
|
||||
make_utf16le_S (w3, w6, w7);
|
||||
make_utf16le_S (w2, w4, w5);
|
||||
make_utf16le_S (w1, w2, w3);
|
||||
make_utf16le_S (w0, w0, w1);
|
||||
make_utf16le (w3, w6, w7);
|
||||
make_utf16le (w2, w4, w5);
|
||||
make_utf16le (w1, w2, w3);
|
||||
make_utf16le (w0, w0, w1);
|
||||
|
||||
sha512_update_vector_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7, 64 * 2);
|
||||
}
|
||||
@ -2421,10 +2421,10 @@ void sha512_update_vector_utf16le (sha512_ctx_vector_t *ctx, const u32x *w, cons
|
||||
w3[2] = w[pos4 + 14];
|
||||
w3[3] = w[pos4 + 15];
|
||||
|
||||
make_utf16le_S (w3, w6, w7);
|
||||
make_utf16le_S (w2, w4, w5);
|
||||
make_utf16le_S (w1, w2, w3);
|
||||
make_utf16le_S (w0, w0, w1);
|
||||
make_utf16le (w3, w6, w7);
|
||||
make_utf16le (w2, w4, w5);
|
||||
make_utf16le (w1, w2, w3);
|
||||
make_utf16le (w0, w0, w1);
|
||||
|
||||
sha512_update_vector_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7, (len - pos1) * 2);
|
||||
}
|
||||
@ -2462,10 +2462,10 @@ void sha512_update_vector_utf16le_swap (sha512_ctx_vector_t *ctx, const u32x *w,
|
||||
w3[2] = w[pos4 + 14];
|
||||
w3[3] = w[pos4 + 15];
|
||||
|
||||
make_utf16le_S (w3, w6, w7);
|
||||
make_utf16le_S (w2, w4, w5);
|
||||
make_utf16le_S (w1, w2, w3);
|
||||
make_utf16le_S (w0, w0, w1);
|
||||
make_utf16le (w3, w6, w7);
|
||||
make_utf16le (w2, w4, w5);
|
||||
make_utf16le (w1, w2, w3);
|
||||
make_utf16le (w0, w0, w1);
|
||||
|
||||
w0[0] = swap32 (w0[0]);
|
||||
w0[1] = swap32 (w0[1]);
|
||||
@ -2520,10 +2520,10 @@ void sha512_update_vector_utf16le_swap (sha512_ctx_vector_t *ctx, const u32x *w,
|
||||
w3[2] = w[pos4 + 14];
|
||||
w3[3] = w[pos4 + 15];
|
||||
|
||||
make_utf16le_S (w3, w6, w7);
|
||||
make_utf16le_S (w2, w4, w5);
|
||||
make_utf16le_S (w1, w2, w3);
|
||||
make_utf16le_S (w0, w0, w1);
|
||||
make_utf16le (w3, w6, w7);
|
||||
make_utf16le (w2, w4, w5);
|
||||
make_utf16le (w1, w2, w3);
|
||||
make_utf16le (w0, w0, w1);
|
||||
|
||||
w0[0] = swap32 (w0[0]);
|
||||
w0[1] = swap32 (w0[1]);
|
||||
@ -2594,10 +2594,10 @@ void sha512_update_vector_utf16beN (sha512_ctx_vector_t *ctx, const u32x *w, con
|
||||
w3[2] = w[pos4 + 14];
|
||||
w3[3] = w[pos4 + 15];
|
||||
|
||||
make_utf16beN_S (w3, w6, w7);
|
||||
make_utf16beN_S (w2, w4, w5);
|
||||
make_utf16beN_S (w1, w2, w3);
|
||||
make_utf16beN_S (w0, w0, w1);
|
||||
make_utf16beN (w3, w6, w7);
|
||||
make_utf16beN (w2, w4, w5);
|
||||
make_utf16beN (w1, w2, w3);
|
||||
make_utf16beN (w0, w0, w1);
|
||||
|
||||
sha512_update_vector_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7, 64 * 2);
|
||||
}
|
||||
@ -2619,10 +2619,10 @@ void sha512_update_vector_utf16beN (sha512_ctx_vector_t *ctx, const u32x *w, con
|
||||
w3[2] = w[pos4 + 14];
|
||||
w3[3] = w[pos4 + 15];
|
||||
|
||||
make_utf16beN_S (w3, w6, w7);
|
||||
make_utf16beN_S (w2, w4, w5);
|
||||
make_utf16beN_S (w1, w2, w3);
|
||||
make_utf16beN_S (w0, w0, w1);
|
||||
make_utf16beN (w3, w6, w7);
|
||||
make_utf16beN (w2, w4, w5);
|
||||
make_utf16beN (w1, w2, w3);
|
||||
make_utf16beN (w0, w0, w1);
|
||||
|
||||
sha512_update_vector_128 (ctx, w0, w1, w2, w3, w4, w5, w6, w7, (len - pos1) * 2);
|
||||
}
|
||||
@ -2879,7 +2879,7 @@ void sha512_hmac_update_vector (sha512_hmac_ctx_vector_t *ctx, const u32x *w, co
|
||||
sha512_update_vector (&ctx->ipad, w, len);
|
||||
}
|
||||
|
||||
void sha512_hmac_final_vector (sha512_hmac_ctx_t *ctx)
|
||||
void sha512_hmac_final_vector (sha512_hmac_ctx_vector_t *ctx)
|
||||
{
|
||||
sha512_final_vector (&ctx->ipad);
|
||||
|
||||
|
@ -72,10 +72,10 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_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);
|
||||
}
|
||||
@ -150,10 +150,10 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_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);
|
||||
}
|
||||
|
@ -90,10 +90,10 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_hmac_final_vector (&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);
|
||||
}
|
||||
@ -188,10 +188,10 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_hmac_final_vector (&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);
|
||||
}
|
||||
|
@ -74,10 +74,10 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_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);
|
||||
}
|
||||
@ -154,10 +154,10 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_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);
|
||||
}
|
||||
|
@ -92,10 +92,10 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_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);
|
||||
}
|
||||
@ -192,10 +192,10 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
md5_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);
|
||||
}
|
||||
|
@ -72,10 +72,10 @@ __kernel void m00150_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);
|
||||
}
|
||||
@ -150,10 +150,10 @@ __kernel void m00150_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);
|
||||
}
|
||||
|
@ -90,10 +90,10 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha1_hmac_final_vector (&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);
|
||||
}
|
||||
@ -188,10 +188,10 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha1_hmac_final_vector (&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);
|
||||
}
|
||||
|
@ -74,10 +74,10 @@ __kernel void m00160_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);
|
||||
}
|
||||
@ -154,10 +154,10 @@ __kernel void m00160_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);
|
||||
}
|
||||
|
@ -92,10 +92,10 @@ __kernel void m00160_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);
|
||||
}
|
||||
@ -192,10 +192,10 @@ __kernel void m00160_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);
|
||||
}
|
||||
|
@ -72,10 +72,10 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha256_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);
|
||||
}
|
||||
@ -150,10 +150,10 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha256_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);
|
||||
}
|
||||
|
@ -90,10 +90,10 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha256_hmac_final_vector (&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);
|
||||
}
|
||||
@ -188,10 +188,10 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha256_hmac_final_vector (&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);
|
||||
}
|
||||
|
@ -74,10 +74,10 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha256_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);
|
||||
}
|
||||
@ -154,10 +154,10 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha256_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);
|
||||
}
|
||||
|
@ -92,10 +92,10 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha256_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);
|
||||
}
|
||||
@ -192,10 +192,10 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
|
||||
sha256_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);
|
||||
}
|
||||
|
Loading…
Reference in New Issue
Block a user