mirror of
https://github.com/hashcat/hashcat.git
synced 2024-12-27 17:08:12 +00:00
Merge remote-tracking branch 'upstream/master'
This commit is contained in:
Royce Williams
commit
176f66086e
29946
OpenCL/inc_common.cl
29946
OpenCL/inc_common.cl
File diff suppressed because it is too large
Load Diff
@ -111,7 +111,7 @@ void md4_init (md4_ctx_t *ctx)
|
||||
void md4_update_64 (md4_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
@ -1047,6 +1047,7 @@ void md4_hmac_update_utf16le_swap (md4_hmac_ctx_t *ctx, const u32 *w, const int
|
||||
{
|
||||
md4_update_utf16le_swap (&ctx->ipad, w, len);
|
||||
}
|
||||
|
||||
void md4_hmac_update_global (md4_hmac_ctx_t *ctx, const __global u32 *w, const int len)
|
||||
{
|
||||
md4_update_global (&ctx->ipad, w, len);
|
||||
@ -1234,7 +1235,7 @@ void md4_init_vector_from_scalar (md4_ctx_vector_t *ctx, md4_ctx_t *ctx0)
|
||||
void md4_update_vector_64 (md4_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
|
||||
@ -145,7 +145,7 @@ void md5_init (md5_ctx_t *ctx)
|
||||
void md5_update_64 (md5_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
@ -1303,7 +1303,7 @@ void md5_init_vector_from_scalar (md5_ctx_vector_t *ctx, md5_ctx_t *ctx0)
|
||||
void md5_update_vector_64 (md5_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
|
||||
@ -245,7 +245,7 @@ void ripemd160_init (ripemd160_ctx_t *ctx)
|
||||
void ripemd160_update_64 (ripemd160_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
@ -1504,7 +1504,7 @@ void ripemd160_init_vector_from_scalar (ripemd160_ctx_vector_t *ctx, ripemd160_c
|
||||
void ripemd160_update_vector_64 (ripemd160_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
|
||||
@ -177,7 +177,7 @@ void sha1_init (sha1_ctx_t *ctx)
|
||||
void sha1_update_64 (sha1_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
@ -1368,7 +1368,7 @@ void sha1_init_vector_from_scalar (sha1_ctx_vector_t *ctx, sha1_ctx_t *ctx0)
|
||||
void sha1_update_vector_64 (sha1_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
|
||||
@ -162,7 +162,7 @@ void sha224_init (sha224_ctx_t *ctx)
|
||||
void sha224_update_64 (sha224_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
@ -1321,7 +1321,7 @@ void sha224_init_vector_from_scalar (sha224_ctx_vector_t *ctx, sha224_ctx_t *ctx
|
||||
void sha224_update_vector_64 (sha224_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
|
||||
@ -162,7 +162,7 @@ void sha256_init (sha256_ctx_t *ctx)
|
||||
void sha256_update_64 (sha256_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
@ -1321,7 +1321,7 @@ void sha256_init_vector_from_scalar (sha256_ctx_vector_t *ctx, sha256_ctx_t *ctx
|
||||
void sha256_update_vector_64 (sha256_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
|
||||
@ -186,7 +186,7 @@ void sha384_init (sha384_ctx_t *ctx)
|
||||
void sha384_update_128 (sha384_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u32 w5[4], u32 w6[4], u32 w7[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 127;
|
||||
const int pos = ctx->len & 127;
|
||||
#else
|
||||
const int pos = ctx->len & 127;
|
||||
#endif
|
||||
@ -2017,7 +2017,7 @@ void sha384_init_vector_from_scalar (sha384_ctx_vector_t *ctx, sha384_ctx_t *ctx
|
||||
void sha384_update_vector_128 (sha384_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x w4[4], u32x w5[4], u32x w6[4], u32x w7[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 127;
|
||||
const int pos = ctx->len & 127;
|
||||
#else
|
||||
const int pos = ctx->len & 127;
|
||||
#endif
|
||||
|
||||
@ -186,7 +186,7 @@ void sha512_init (sha512_ctx_t *ctx)
|
||||
void sha512_update_128 (sha512_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], u32 w4[4], u32 w5[4], u32 w6[4], u32 w7[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 127;
|
||||
const int pos = ctx->len & 127;
|
||||
#else
|
||||
const int pos = ctx->len & 127;
|
||||
#endif
|
||||
@ -2017,7 +2017,7 @@ void sha512_init_vector_from_scalar (sha512_ctx_vector_t *ctx, sha512_ctx_t *ctx
|
||||
void sha512_update_vector_128 (sha512_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x w4[4], u32x w5[4], u32x w6[4], u32x w7[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 127;
|
||||
const int pos = ctx->len & 127;
|
||||
#else
|
||||
const int pos = ctx->len & 127;
|
||||
#endif
|
||||
|
||||
@ -1345,7 +1345,7 @@ void whirlpool_init (whirlpool_ctx_t *ctx, __local u32 (*s_Ch)[256], __local u32
|
||||
void whirlpool_update_64 (whirlpool_ctx_t *ctx, u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
@ -2608,7 +2608,7 @@ void whirlpool_init_vector_from_scalar (whirlpool_ctx_vector_t *ctx, whirlpool_c
|
||||
void whirlpool_update_vector_64 (whirlpool_ctx_vector_t *ctx, u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], const int len)
|
||||
{
|
||||
#ifdef IS_AMD
|
||||
volatile const int pos = ctx->len & 63;
|
||||
const int pos = ctx->len & 63;
|
||||
#else
|
||||
const int pos = ctx->len & 63;
|
||||
#endif
|
||||
|
||||
324
OpenCL/inc_rp.cl
324
OpenCL/inc_rp.cl
@ -3,7 +3,7 @@
|
||||
* License.....: MIT
|
||||
*/
|
||||
|
||||
inline u32 generate_cmask (const u32 value)
|
||||
u32 generate_cmask (const u32 value)
|
||||
{
|
||||
const u32 rmask = ((value & 0x40404040u) >> 1u)
|
||||
& ~((value & 0x80808080u) >> 2u);
|
||||
@ -14,7 +14,7 @@ inline u32 generate_cmask (const u32 value)
|
||||
return rmask & ~hmask & lmask;
|
||||
}
|
||||
|
||||
inline void truncate_right (u32 buf0[4], u32 buf1[4], const u32 offset)
|
||||
void truncate_right (u32 buf0[4], u32 buf1[4], const u32 offset)
|
||||
{
|
||||
const u32 tmp = (1u << ((offset & 3u) * 8u)) - 1u;
|
||||
|
||||
@ -67,7 +67,7 @@ inline void truncate_right (u32 buf0[4], u32 buf1[4], const u32 offset)
|
||||
}
|
||||
}
|
||||
|
||||
inline void truncate_left (u32 buf0[4], u32 buf1[4], const u32 offset)
|
||||
void truncate_left (u32 buf0[4], u32 buf1[4], const u32 offset)
|
||||
{
|
||||
const u32 tmp = ~((1u << ((offset & 3u) * 8u)) - 1u);
|
||||
|
||||
@ -120,7 +120,7 @@ inline void truncate_left (u32 buf0[4], u32 buf1[4], const u32 offset)
|
||||
}
|
||||
}
|
||||
|
||||
inline void lshift_block (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 out1[4])
|
||||
void lshift_block (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 out1[4])
|
||||
{
|
||||
out0[0] = amd_bytealign_S (in0[1], in0[0], 1);
|
||||
out0[1] = amd_bytealign_S (in0[2], in0[1], 1);
|
||||
@ -132,7 +132,7 @@ inline void lshift_block (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 o
|
||||
out1[3] = amd_bytealign_S ( 0, in1[3], 1);
|
||||
}
|
||||
|
||||
inline void rshift_block (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 out1[4])
|
||||
void rshift_block (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 out1[4])
|
||||
{
|
||||
out1[3] = amd_bytealign_S (in1[3], in1[2], 3);
|
||||
out1[2] = amd_bytealign_S (in1[2], in1[1], 3);
|
||||
@ -144,7 +144,7 @@ inline void rshift_block (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 o
|
||||
out0[0] = amd_bytealign_S (in0[0], 0, 3);
|
||||
}
|
||||
|
||||
inline void lshift_block_N (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 out1[4], const u32 num)
|
||||
void lshift_block_N (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 out1[4], const u32 num)
|
||||
{
|
||||
switch (num)
|
||||
{
|
||||
@ -439,7 +439,7 @@ inline void lshift_block_N (const u32 in0[4], const u32 in1[4], u32 out0[4], u32
|
||||
}
|
||||
}
|
||||
|
||||
inline void rshift_block_N (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 out1[4], const u32 num)
|
||||
void rshift_block_N (const u32 in0[4], const u32 in1[4], u32 out0[4], u32 out1[4], const u32 num)
|
||||
{
|
||||
switch (num)
|
||||
{
|
||||
@ -734,7 +734,7 @@ inline void rshift_block_N (const u32 in0[4], const u32 in1[4], u32 out0[4], u32
|
||||
}
|
||||
}
|
||||
|
||||
inline void append_block1 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32 src_r0)
|
||||
void append_block1 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32 src_r0)
|
||||
{
|
||||
// this version works with 1 byte append only
|
||||
|
||||
@ -754,12 +754,11 @@ inline void append_block1 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
buf1[3] |= (offset >= 28) ? tmp : 0;
|
||||
}
|
||||
|
||||
inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32 src_l0[4], const u32 src_l1[4], const u32 src_r0[4], const u32 src_r1[4])
|
||||
void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32 src_l0[4], const u32 src_l1[4], const u32 src_r0[4], const u32 src_r1[4])
|
||||
{
|
||||
#if defined IS_AMD || defined IS_GENERIC
|
||||
const int offset_mod_4 = offset & 3;
|
||||
|
||||
const int offset_minus_4 = 4 - offset;
|
||||
const int offset_minus_4 = 4 - offset_mod_4;
|
||||
|
||||
u32 s0 = 0;
|
||||
u32 s1 = 0;
|
||||
@ -769,64 +768,69 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
u32 s5 = 0;
|
||||
u32 s6 = 0;
|
||||
u32 s7 = 0;
|
||||
u32 s8 = 0;
|
||||
|
||||
#if defined IS_AMD || defined IS_GENERIC
|
||||
const u32 src_r00 = swap32_S (src_r0[0]);
|
||||
const u32 src_r01 = swap32_S (src_r0[1]);
|
||||
const u32 src_r02 = swap32_S (src_r0[2]);
|
||||
const u32 src_r03 = swap32_S (src_r0[3]);
|
||||
const u32 src_r10 = swap32_S (src_r1[0]);
|
||||
const u32 src_r11 = swap32_S (src_r1[1]);
|
||||
const u32 src_r12 = swap32_S (src_r1[2]);
|
||||
const u32 src_r13 = swap32_S (src_r1[3]);
|
||||
|
||||
switch (offset / 4)
|
||||
{
|
||||
case 0:
|
||||
s8 = amd_bytealign_S ( 0, src_r1[3], offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r1[3], src_r1[2], offset_minus_4);
|
||||
s6 = amd_bytealign_S (src_r1[2], src_r1[1], offset_minus_4);
|
||||
s5 = amd_bytealign_S (src_r1[1], src_r1[0], offset_minus_4);
|
||||
s4 = amd_bytealign_S (src_r1[0], src_r0[3], offset_minus_4);
|
||||
s3 = amd_bytealign_S (src_r0[3], src_r0[2], offset_minus_4);
|
||||
s2 = amd_bytealign_S (src_r0[2], src_r0[1], offset_minus_4);
|
||||
s1 = amd_bytealign_S (src_r0[1], src_r0[0], offset_minus_4);
|
||||
s0 = amd_bytealign_S (src_r0[0], 0, offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r12, src_r13, offset);
|
||||
s6 = amd_bytealign_S (src_r11, src_r12, offset);
|
||||
s5 = amd_bytealign_S (src_r10, src_r11, offset);
|
||||
s4 = amd_bytealign_S (src_r03, src_r10, offset);
|
||||
s3 = amd_bytealign_S (src_r02, src_r03, offset);
|
||||
s2 = amd_bytealign_S (src_r01, src_r02, offset);
|
||||
s1 = amd_bytealign_S (src_r00, src_r01, offset);
|
||||
s0 = amd_bytealign_S ( 0, src_r00, offset);
|
||||
break;
|
||||
|
||||
case 1:
|
||||
s8 = amd_bytealign_S ( 0, src_r1[2], offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r1[2], src_r1[1], offset_minus_4);
|
||||
s6 = amd_bytealign_S (src_r1[1], src_r1[0], offset_minus_4);
|
||||
s5 = amd_bytealign_S (src_r1[0], src_r0[3], offset_minus_4);
|
||||
s4 = amd_bytealign_S (src_r0[3], src_r0[2], offset_minus_4);
|
||||
s3 = amd_bytealign_S (src_r0[2], src_r0[1], offset_minus_4);
|
||||
s2 = amd_bytealign_S (src_r0[1], src_r0[0], offset_minus_4);
|
||||
s1 = amd_bytealign_S (src_r0[0], 0, offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r11, src_r12, offset);
|
||||
s6 = amd_bytealign_S (src_r10, src_r11, offset);
|
||||
s5 = amd_bytealign_S (src_r03, src_r10, offset);
|
||||
s4 = amd_bytealign_S (src_r02, src_r03, offset);
|
||||
s3 = amd_bytealign_S (src_r01, src_r02, offset);
|
||||
s2 = amd_bytealign_S (src_r00, src_r01, offset);
|
||||
s1 = amd_bytealign_S ( 0, src_r00, offset);
|
||||
s0 = 0;
|
||||
break;
|
||||
|
||||
case 2:
|
||||
s8 = amd_bytealign_S ( 0, src_r1[1], offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r1[1], src_r1[0], offset_minus_4);
|
||||
s6 = amd_bytealign_S (src_r1[0], src_r0[3], offset_minus_4);
|
||||
s5 = amd_bytealign_S (src_r0[3], src_r0[2], offset_minus_4);
|
||||
s4 = amd_bytealign_S (src_r0[2], src_r0[1], offset_minus_4);
|
||||
s3 = amd_bytealign_S (src_r0[1], src_r0[0], offset_minus_4);
|
||||
s2 = amd_bytealign_S (src_r0[0], 0, offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r10, src_r11, offset);
|
||||
s6 = amd_bytealign_S (src_r03, src_r10, offset);
|
||||
s5 = amd_bytealign_S (src_r02, src_r03, offset);
|
||||
s4 = amd_bytealign_S (src_r01, src_r02, offset);
|
||||
s3 = amd_bytealign_S (src_r00, src_r01, offset);
|
||||
s2 = amd_bytealign_S ( 0, src_r00, offset);
|
||||
s1 = 0;
|
||||
s0 = 0;
|
||||
break;
|
||||
|
||||
case 3:
|
||||
s8 = amd_bytealign_S ( 0, src_r1[0], offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r1[0], src_r0[3], offset_minus_4);
|
||||
s6 = amd_bytealign_S (src_r0[3], src_r0[2], offset_minus_4);
|
||||
s5 = amd_bytealign_S (src_r0[2], src_r0[1], offset_minus_4);
|
||||
s4 = amd_bytealign_S (src_r0[1], src_r0[0], offset_minus_4);
|
||||
s3 = amd_bytealign_S (src_r0[0], 0, offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r03, src_r10, offset);
|
||||
s6 = amd_bytealign_S (src_r02, src_r03, offset);
|
||||
s5 = amd_bytealign_S (src_r01, src_r02, offset);
|
||||
s4 = amd_bytealign_S (src_r00, src_r01, offset);
|
||||
s3 = amd_bytealign_S ( 0, src_r00, offset);
|
||||
s2 = 0;
|
||||
s1 = 0;
|
||||
s0 = 0;
|
||||
|
||||
break;
|
||||
|
||||
case 4:
|
||||
s8 = amd_bytealign_S ( 0, src_r0[3], offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r0[3], src_r0[2], offset_minus_4);
|
||||
s6 = amd_bytealign_S (src_r0[2], src_r0[1], offset_minus_4);
|
||||
s5 = amd_bytealign_S (src_r0[1], src_r0[0], offset_minus_4);
|
||||
s4 = amd_bytealign_S (src_r0[0], 0, offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r02, src_r03, offset);
|
||||
s6 = amd_bytealign_S (src_r01, src_r02, offset);
|
||||
s5 = amd_bytealign_S (src_r00, src_r01, offset);
|
||||
s4 = amd_bytealign_S ( 0, src_r00, offset);
|
||||
s3 = 0;
|
||||
s2 = 0;
|
||||
s1 = 0;
|
||||
@ -834,10 +838,9 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
break;
|
||||
|
||||
case 5:
|
||||
s8 = amd_bytealign_S ( 0, src_r0[2], offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r0[2], src_r0[1], offset_minus_4);
|
||||
s6 = amd_bytealign_S (src_r0[1], src_r0[0], offset_minus_4);
|
||||
s5 = amd_bytealign_S (src_r0[0], 0, offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r01, src_r02, offset);
|
||||
s6 = amd_bytealign_S (src_r00, src_r01, offset);
|
||||
s5 = amd_bytealign_S ( 0, src_r00, offset);
|
||||
s4 = 0;
|
||||
s3 = 0;
|
||||
s2 = 0;
|
||||
@ -846,9 +849,8 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
break;
|
||||
|
||||
case 6:
|
||||
s8 = amd_bytealign_S ( 0, src_r0[1], offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r0[1], src_r0[0], offset_minus_4);
|
||||
s6 = amd_bytealign_S (src_r0[0], 0, offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r00, src_r01, offset);
|
||||
s6 = amd_bytealign_S ( 0, src_r00, offset);
|
||||
s5 = 0;
|
||||
s4 = 0;
|
||||
s3 = 0;
|
||||
@ -858,8 +860,7 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
break;
|
||||
|
||||
case 7:
|
||||
s8 = amd_bytealign_S ( 0, src_r0[0], offset_minus_4);
|
||||
s7 = amd_bytealign_S (src_r0[0], 0, offset_minus_4);
|
||||
s7 = amd_bytealign_S ( 0, src_r00, offset);
|
||||
s6 = 0;
|
||||
s5 = 0;
|
||||
s4 = 0;
|
||||
@ -870,85 +871,69 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
break;
|
||||
}
|
||||
|
||||
if (offset_mod_4 == 0)
|
||||
{
|
||||
buf0[0] = src_l0[0] | s1;
|
||||
buf0[1] = src_l0[1] | s2;
|
||||
buf0[2] = src_l0[2] | s3;
|
||||
buf0[3] = src_l0[3] | s4;
|
||||
buf1[0] = src_l1[0] | s5;
|
||||
buf1[1] = src_l1[1] | s6;
|
||||
buf1[2] = src_l1[2] | s7;
|
||||
buf1[3] = src_l1[3] | s8;
|
||||
}
|
||||
else
|
||||
{
|
||||
buf0[0] = src_l0[0] | s0;
|
||||
buf0[1] = src_l0[1] | s1;
|
||||
buf0[2] = src_l0[2] | s2;
|
||||
buf0[3] = src_l0[3] | s3;
|
||||
buf1[0] = src_l1[0] | s4;
|
||||
buf1[1] = src_l1[1] | s5;
|
||||
buf1[2] = src_l1[2] | s6;
|
||||
buf1[3] = src_l1[3] | s7;
|
||||
}
|
||||
s0 = swap32_S (s0);
|
||||
s1 = swap32_S (s1);
|
||||
s2 = swap32_S (s2);
|
||||
s3 = swap32_S (s3);
|
||||
s4 = swap32_S (s4);
|
||||
s5 = swap32_S (s5);
|
||||
s6 = swap32_S (s6);
|
||||
s7 = swap32_S (s7);
|
||||
#endif
|
||||
|
||||
#ifdef IS_NV
|
||||
const int offset_minus_4 = 4 - (offset % 4);
|
||||
|
||||
const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
|
||||
|
||||
u32 s0 = 0;
|
||||
u32 s1 = 0;
|
||||
u32 s2 = 0;
|
||||
u32 s3 = 0;
|
||||
u32 s4 = 0;
|
||||
u32 s5 = 0;
|
||||
u32 s6 = 0;
|
||||
u32 s7 = 0;
|
||||
const u32 src_r00 = src_r0[0];
|
||||
const u32 src_r01 = src_r0[1];
|
||||
const u32 src_r02 = src_r0[2];
|
||||
const u32 src_r03 = src_r0[3];
|
||||
const u32 src_r10 = src_r1[0];
|
||||
const u32 src_r11 = src_r1[1];
|
||||
const u32 src_r12 = src_r1[2];
|
||||
const u32 src_r13 = src_r1[3];
|
||||
|
||||
switch (offset / 4)
|
||||
{
|
||||
case 0:
|
||||
s7 = __byte_perm_S (src_r1[2], src_r1[3], selector);
|
||||
s6 = __byte_perm_S (src_r1[1], src_r1[2], selector);
|
||||
s5 = __byte_perm_S (src_r1[0], src_r1[1], selector);
|
||||
s4 = __byte_perm_S (src_r0[3], src_r1[0], selector);
|
||||
s3 = __byte_perm_S (src_r0[2], src_r0[3], selector);
|
||||
s2 = __byte_perm_S (src_r0[1], src_r0[2], selector);
|
||||
s1 = __byte_perm_S (src_r0[0], src_r0[1], selector);
|
||||
s0 = __byte_perm_S ( 0, src_r0[0], selector);
|
||||
s7 = __byte_perm_S (src_r12, src_r13, selector);
|
||||
s6 = __byte_perm_S (src_r11, src_r12, selector);
|
||||
s5 = __byte_perm_S (src_r10, src_r11, selector);
|
||||
s4 = __byte_perm_S (src_r03, src_r10, selector);
|
||||
s3 = __byte_perm_S (src_r02, src_r03, selector);
|
||||
s2 = __byte_perm_S (src_r01, src_r02, selector);
|
||||
s1 = __byte_perm_S (src_r00, src_r01, selector);
|
||||
s0 = __byte_perm_S ( 0, src_r00, selector);
|
||||
break;
|
||||
|
||||
case 1:
|
||||
s7 = __byte_perm_S (src_r1[1], src_r1[2], selector);
|
||||
s6 = __byte_perm_S (src_r1[0], src_r1[1], selector);
|
||||
s5 = __byte_perm_S (src_r0[3], src_r1[0], selector);
|
||||
s4 = __byte_perm_S (src_r0[2], src_r0[3], selector);
|
||||
s3 = __byte_perm_S (src_r0[1], src_r0[2], selector);
|
||||
s2 = __byte_perm_S (src_r0[0], src_r0[1], selector);
|
||||
s1 = __byte_perm_S ( 0, src_r0[0], selector);
|
||||
s7 = __byte_perm_S (src_r11, src_r12, selector);
|
||||
s6 = __byte_perm_S (src_r10, src_r11, selector);
|
||||
s5 = __byte_perm_S (src_r03, src_r10, selector);
|
||||
s4 = __byte_perm_S (src_r02, src_r03, selector);
|
||||
s3 = __byte_perm_S (src_r01, src_r02, selector);
|
||||
s2 = __byte_perm_S (src_r00, src_r01, selector);
|
||||
s1 = __byte_perm_S ( 0, src_r00, selector);
|
||||
s0 = 0;
|
||||
break;
|
||||
|
||||
case 2:
|
||||
s7 = __byte_perm_S (src_r1[0], src_r1[1], selector);
|
||||
s6 = __byte_perm_S (src_r0[3], src_r1[0], selector);
|
||||
s5 = __byte_perm_S (src_r0[2], src_r0[3], selector);
|
||||
s4 = __byte_perm_S (src_r0[1], src_r0[2], selector);
|
||||
s3 = __byte_perm_S (src_r0[0], src_r0[1], selector);
|
||||
s2 = __byte_perm_S ( 0, src_r0[0], selector);
|
||||
s7 = __byte_perm_S (src_r10, src_r11, selector);
|
||||
s6 = __byte_perm_S (src_r03, src_r10, selector);
|
||||
s5 = __byte_perm_S (src_r02, src_r03, selector);
|
||||
s4 = __byte_perm_S (src_r01, src_r02, selector);
|
||||
s3 = __byte_perm_S (src_r00, src_r01, selector);
|
||||
s2 = __byte_perm_S ( 0, src_r00, selector);
|
||||
s1 = 0;
|
||||
s0 = 0;
|
||||
break;
|
||||
|
||||
case 3:
|
||||
s7 = __byte_perm_S (src_r0[3], src_r1[0], selector);
|
||||
s6 = __byte_perm_S (src_r0[2], src_r0[3], selector);
|
||||
s5 = __byte_perm_S (src_r0[1], src_r0[2], selector);
|
||||
s4 = __byte_perm_S (src_r0[0], src_r0[1], selector);
|
||||
s3 = __byte_perm_S ( 0, src_r0[0], selector);
|
||||
s7 = __byte_perm_S (src_r03, src_r10, selector);
|
||||
s6 = __byte_perm_S (src_r02, src_r03, selector);
|
||||
s5 = __byte_perm_S (src_r01, src_r02, selector);
|
||||
s4 = __byte_perm_S (src_r00, src_r01, selector);
|
||||
s3 = __byte_perm_S ( 0, src_r00, selector);
|
||||
s2 = 0;
|
||||
s1 = 0;
|
||||
s0 = 0;
|
||||
@ -956,10 +941,10 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
break;
|
||||
|
||||
case 4:
|
||||
s7 = __byte_perm_S (src_r0[2], src_r0[3], selector);
|
||||
s6 = __byte_perm_S (src_r0[1], src_r0[2], selector);
|
||||
s5 = __byte_perm_S (src_r0[0], src_r0[1], selector);
|
||||
s4 = __byte_perm_S ( 0, src_r0[0], selector);
|
||||
s7 = __byte_perm_S (src_r02, src_r03, selector);
|
||||
s6 = __byte_perm_S (src_r01, src_r02, selector);
|
||||
s5 = __byte_perm_S (src_r00, src_r01, selector);
|
||||
s4 = __byte_perm_S ( 0, src_r00, selector);
|
||||
s3 = 0;
|
||||
s2 = 0;
|
||||
s1 = 0;
|
||||
@ -967,9 +952,9 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
break;
|
||||
|
||||
case 5:
|
||||
s7 = __byte_perm_S (src_r0[1], src_r0[2], selector);
|
||||
s6 = __byte_perm_S (src_r0[0], src_r0[1], selector);
|
||||
s5 = __byte_perm_S ( 0, src_r0[0], selector);
|
||||
s7 = __byte_perm_S (src_r01, src_r02, selector);
|
||||
s6 = __byte_perm_S (src_r00, src_r01, selector);
|
||||
s5 = __byte_perm_S ( 0, src_r00, selector);
|
||||
s4 = 0;
|
||||
s3 = 0;
|
||||
s2 = 0;
|
||||
@ -978,8 +963,8 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
break;
|
||||
|
||||
case 6:
|
||||
s7 = __byte_perm_S (src_r0[0], src_r0[1], selector);
|
||||
s6 = __byte_perm_S ( 0, src_r0[0], selector);
|
||||
s7 = __byte_perm_S (src_r00, src_r01, selector);
|
||||
s6 = __byte_perm_S ( 0, src_r00, selector);
|
||||
s5 = 0;
|
||||
s4 = 0;
|
||||
s3 = 0;
|
||||
@ -989,7 +974,7 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
break;
|
||||
|
||||
case 7:
|
||||
s7 = __byte_perm_S ( 0, src_r0[0], selector);
|
||||
s7 = __byte_perm_S ( 0, src_r00, selector);
|
||||
s6 = 0;
|
||||
s5 = 0;
|
||||
s4 = 0;
|
||||
@ -999,6 +984,7 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
s0 = 0;
|
||||
break;
|
||||
}
|
||||
#endif
|
||||
|
||||
buf0[0] = src_l0[0] | s0;
|
||||
buf0[1] = src_l0[1] | s1;
|
||||
@ -1008,11 +994,9 @@ inline void append_block8 (const u32 offset, u32 buf0[4], u32 buf1[4], const u32
|
||||
buf1[1] = src_l1[1] | s5;
|
||||
buf1[2] = src_l1[2] | s6;
|
||||
buf1[3] = src_l1[3] | s7;
|
||||
|
||||
#endif
|
||||
}
|
||||
|
||||
inline void reverse_block (u32 in0[4], u32 in1[4], u32 out0[4], u32 out1[4], const u32 len)
|
||||
void reverse_block (u32 in0[4], u32 in1[4], u32 out0[4], u32 out1[4], const u32 len)
|
||||
{
|
||||
rshift_block_N (in0, in1, out0, out1, 32 - len);
|
||||
|
||||
@ -1038,7 +1022,7 @@ inline void reverse_block (u32 in0[4], u32 in1[4], u32 out0[4], u32 out1[4], con
|
||||
out1[3] = swap32_S (tib41[3]);
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_lrest (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_lrest (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
buf0[0] |= (generate_cmask (buf0[0]));
|
||||
buf0[1] |= (generate_cmask (buf0[1]));
|
||||
@ -1052,7 +1036,7 @@ inline u32 rule_op_mangle_lrest (const u32 p0, const u32 p1, u32 buf0[4], u32 bu
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_urest (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_urest (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
buf0[0] &= ~(generate_cmask (buf0[0]));
|
||||
buf0[1] &= ~(generate_cmask (buf0[1]));
|
||||
@ -1066,7 +1050,7 @@ inline u32 rule_op_mangle_urest (const u32 p0, const u32 p1, u32 buf0[4], u32 bu
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_lrest_ufirst (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_lrest_ufirst (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
rule_op_mangle_lrest (p0, p1, buf0, buf1, in_len);
|
||||
|
||||
@ -1075,7 +1059,7 @@ inline u32 rule_op_mangle_lrest_ufirst (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_urest_lfirst (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_urest_lfirst (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
rule_op_mangle_urest (p0, p1, buf0, buf1, in_len);
|
||||
|
||||
@ -1084,7 +1068,7 @@ inline u32 rule_op_mangle_urest_lfirst (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_trest (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_trest (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
buf0[0] ^= (generate_cmask (buf0[0]));
|
||||
buf0[1] ^= (generate_cmask (buf0[1]));
|
||||
@ -1098,7 +1082,7 @@ inline u32 rule_op_mangle_trest (const u32 p0, const u32 p1, u32 buf0[4], u32 bu
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_toggle_at (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_toggle_at (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -1119,14 +1103,14 @@ inline u32 rule_op_mangle_toggle_at (const u32 p0, const u32 p1, u32 buf0[4], u3
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_reverse (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_reverse (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
reverse_block (buf0, buf1, buf0, buf1, in_len);
|
||||
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_dupeword (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_dupeword (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if ((in_len + in_len) >= 32) return (in_len);
|
||||
|
||||
@ -1139,7 +1123,7 @@ inline u32 rule_op_mangle_dupeword (const u32 p0, const u32 p1, u32 buf0[4], u32
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_dupeword_times (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_dupeword_times (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (((in_len * p0) + in_len) >= 32) return (in_len);
|
||||
|
||||
@ -1167,7 +1151,7 @@ inline u32 rule_op_mangle_dupeword_times (const u32 p0, const u32 p1, u32 buf0[4
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_reflect (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_reflect (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if ((in_len + in_len) >= 32) return (in_len);
|
||||
|
||||
@ -1185,7 +1169,7 @@ inline u32 rule_op_mangle_reflect (const u32 p0, const u32 p1, u32 buf0[4], u32
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_append (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_append (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if ((in_len + 1) >= 32) return (in_len);
|
||||
|
||||
@ -1198,7 +1182,7 @@ inline u32 rule_op_mangle_append (const u32 p0, const u32 p1, u32 buf0[4], u32 b
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_prepend (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_prepend (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if ((in_len + 1) >= 32) return (in_len);
|
||||
|
||||
@ -1213,7 +1197,7 @@ inline u32 rule_op_mangle_prepend (const u32 p0, const u32 p1, u32 buf0[4], u32
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_rotate_left (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_rotate_left (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (in_len == 0) return (in_len);
|
||||
|
||||
@ -1237,7 +1221,7 @@ inline u32 rule_op_mangle_rotate_left (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_rotate_right (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_rotate_right (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (in_len == 0) return (in_len);
|
||||
|
||||
@ -1267,7 +1251,7 @@ inline u32 rule_op_mangle_rotate_right (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_delete_first (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_delete_first (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (in_len == 0) return (in_len);
|
||||
|
||||
@ -1278,7 +1262,7 @@ inline u32 rule_op_mangle_delete_first (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len1;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_delete_last (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_delete_last (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (in_len == 0) return (in_len);
|
||||
|
||||
@ -1298,7 +1282,7 @@ inline u32 rule_op_mangle_delete_last (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len1;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_delete_at (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_delete_at (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -1373,7 +1357,7 @@ inline u32 rule_op_mangle_delete_at (const u32 p0, const u32 p1, u32 buf0[4], u3
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_extract (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_extract (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -1388,7 +1372,7 @@ inline u32 rule_op_mangle_extract (const u32 p0, const u32 p1, u32 buf0[4], u32
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_omit (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_omit (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -1474,7 +1458,7 @@ inline u32 rule_op_mangle_omit (const u32 p0, const u32 p1, u32 buf0[4], u32 buf
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_insert (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_insert (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 > in_len) return (in_len);
|
||||
|
||||
@ -1546,7 +1530,7 @@ inline u32 rule_op_mangle_insert (const u32 p0, const u32 p1, u32 buf0[4], u32 b
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_overstrike (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_overstrike (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -1569,7 +1553,7 @@ inline u32 rule_op_mangle_overstrike (const u32 p0, const u32 p1, u32 buf0[4], u
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_truncate_at (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_truncate_at (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -1578,7 +1562,7 @@ inline u32 rule_op_mangle_truncate_at (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return p0;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_replace (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_replace (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
const uchar4 tmp0 = (uchar4) (p0);
|
||||
const uchar4 tmp1 = (uchar4) (p1);
|
||||
@ -1597,7 +1581,7 @@ inline u32 rule_op_mangle_replace (const u32 p0, const u32 p1, u32 buf0[4], u32
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_purgechar (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_purgechar (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
u32 out_len = 0;
|
||||
|
||||
@ -1638,13 +1622,13 @@ inline u32 rule_op_mangle_purgechar (const u32 p0, const u32 p1, u32 buf0[4], u3
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_togglecase_rec (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_togglecase_rec (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
// TODO
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_dupechar_first (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_dupechar_first (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if ( in_len == 0) return (in_len);
|
||||
if ((in_len + p0) >= 32) return (in_len);
|
||||
@ -1831,7 +1815,7 @@ inline u32 rule_op_mangle_dupechar_first (const u32 p0, const u32 p1, u32 buf0[4
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_dupechar_last (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_dupechar_last (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if ( in_len == 0) return (in_len);
|
||||
if ((in_len + p0) >= 32) return (in_len);
|
||||
@ -1865,7 +1849,7 @@ inline u32 rule_op_mangle_dupechar_last (const u32 p0, const u32 p1, u32 buf0[4]
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_dupechar_all (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_dupechar_all (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if ( in_len == 0) return (in_len);
|
||||
if ((in_len + in_len) >= 32) return (in_len);
|
||||
@ -1898,7 +1882,7 @@ inline u32 rule_op_mangle_dupechar_all (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_switch_first (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_switch_first (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (in_len < 2) return (in_len);
|
||||
|
||||
@ -1907,7 +1891,7 @@ inline u32 rule_op_mangle_switch_first (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_switch_last (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_switch_last (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (in_len < 2) return (in_len);
|
||||
|
||||
@ -1992,7 +1976,7 @@ inline u32 rule_op_mangle_switch_last (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_switch_at (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_switch_at (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
if (p1 >= in_len) return (in_len);
|
||||
@ -2239,7 +2223,7 @@ inline u32 rule_op_mangle_switch_at (const u32 p0, const u32 p1, u32 buf0[4], u3
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_chr_shiftl (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_chr_shiftl (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -2261,7 +2245,7 @@ inline u32 rule_op_mangle_chr_shiftl (const u32 p0, const u32 p1, u32 buf0[4], u
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_chr_shiftr (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_chr_shiftr (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -2283,7 +2267,7 @@ inline u32 rule_op_mangle_chr_shiftr (const u32 p0, const u32 p1, u32 buf0[4], u
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_chr_incr (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_chr_incr (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -2307,7 +2291,7 @@ inline u32 rule_op_mangle_chr_incr (const u32 p0, const u32 p1, u32 buf0[4], u32
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_chr_decr (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_chr_decr (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 >= in_len) return (in_len);
|
||||
|
||||
@ -2331,7 +2315,7 @@ inline u32 rule_op_mangle_chr_decr (const u32 p0, const u32 p1, u32 buf0[4], u32
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_replace_np1 (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_replace_np1 (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if ((p0 + 1) >= in_len) return (in_len);
|
||||
|
||||
@ -2358,7 +2342,7 @@ inline u32 rule_op_mangle_replace_np1 (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_replace_nm1 (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_replace_nm1 (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 == 0) return (in_len);
|
||||
|
||||
@ -2387,7 +2371,7 @@ inline u32 rule_op_mangle_replace_nm1 (const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_dupeblock_first (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_dupeblock_first (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 > in_len) return (in_len);
|
||||
|
||||
@ -2425,7 +2409,7 @@ inline u32 rule_op_mangle_dupeblock_first (const u32 p0, const u32 p1, u32 buf0[
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_dupeblock_last (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_dupeblock_last (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
if (p0 > in_len) return (in_len);
|
||||
|
||||
@ -2454,7 +2438,7 @@ inline u32 rule_op_mangle_dupeblock_last (const u32 p0, const u32 p1, u32 buf0[4
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 rule_op_mangle_title_sep (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 rule_op_mangle_title_sep (const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
buf0[0] |= (generate_cmask (buf0[0]));
|
||||
buf0[1] |= (generate_cmask (buf0[1]));
|
||||
@ -2497,7 +2481,7 @@ inline u32 rule_op_mangle_title_sep (const u32 p0, const u32 p1, u32 buf0[4], u3
|
||||
return in_len;
|
||||
}
|
||||
|
||||
inline u32 apply_rule (const u32 name, const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
u32 apply_rule (const u32 name, const u32 p0, const u32 p1, u32 buf0[4], u32 buf1[4], const u32 in_len)
|
||||
{
|
||||
u32 out_len = in_len;
|
||||
|
||||
@ -2549,7 +2533,7 @@ inline u32 apply_rule (const u32 name, const u32 p0, const u32 p1, u32 buf0[4],
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32 apply_rules (__global const u32 *cmds, u32 buf0[4], u32 buf1[4], const u32 len)
|
||||
u32 apply_rules (__global const u32 *cmds, u32 buf0[4], u32 buf1[4], const u32 len)
|
||||
{
|
||||
u32 out_len = len;
|
||||
|
||||
@ -2567,7 +2551,7 @@ inline u32 apply_rules (__global const u32 *cmds, u32 buf0[4], u32 buf1[4], cons
|
||||
return out_len;
|
||||
}
|
||||
|
||||
inline u32x apply_rules_vect (const u32 pw_buf0[4], const u32 pw_buf1[4], const u32 pw_len, __global const kernel_rule_t *rules_buf, const u32 il_pos, u32x buf0[4], u32x buf1[4])
|
||||
u32x apply_rules_vect (const u32 pw_buf0[4], const u32 pw_buf1[4], const u32 pw_len, __global const kernel_rule_t *rules_buf, const u32 il_pos, u32x buf0[4], u32x buf1[4])
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
|
||||
|
||||
@ -1054,7 +1054,7 @@
|
||||
|
||||
// attack-mode 0
|
||||
|
||||
inline u32x ix_create_bft (__global const bf_t *bfs_buf, const u32 il_pos)
|
||||
u32x ix_create_bft (__global const bf_t *bfs_buf, const u32 il_pos)
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
const u32x ix = (u32x) (bfs_buf[il_pos + 0].i);
|
||||
@ -1073,7 +1073,7 @@ inline u32x ix_create_bft (__global const bf_t *bfs_buf, const u32 il_pos)
|
||||
|
||||
// attack-mode 1
|
||||
|
||||
inline u32x pwlenx_create_combt (__global const pw_t *combs_buf, const u32 il_pos)
|
||||
u32x pwlenx_create_combt (__global const pw_t *combs_buf, const u32 il_pos)
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
const u32x pw_lenx = (u32x) (combs_buf[il_pos + 0].pw_len);
|
||||
@ -1090,7 +1090,7 @@ inline u32x pwlenx_create_combt (__global const pw_t *combs_buf, const u32 il_po
|
||||
return pw_lenx;
|
||||
}
|
||||
|
||||
inline u32x ix_create_combt (__global const pw_t *combs_buf, const u32 il_pos, const int idx)
|
||||
u32x ix_create_combt (__global const pw_t *combs_buf, const u32 il_pos, const int idx)
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
const u32x ix = (u32x) (combs_buf[il_pos + 0].i[idx]);
|
||||
|
||||
@ -150,7 +150,7 @@ void twofish256_decrypt_xts_next (const u32 *in, u32 *out, u32 *T, u32 *sk, u32
|
||||
|
||||
// 512 bit
|
||||
|
||||
int verify_header_aes (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
int verify_header_aes (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
{
|
||||
u32 ks_aes[60];
|
||||
|
||||
@ -206,7 +206,7 @@ int verify_header_aes (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *u
|
||||
return 1;
|
||||
}
|
||||
|
||||
int verify_header_serpent (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2)
|
||||
int verify_header_serpent (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2)
|
||||
{
|
||||
u32 ks_serpent[140];
|
||||
|
||||
@ -262,7 +262,7 @@ int verify_header_serpent (__global tc_t *esalt_bufs, const u32 *ukey1, const u3
|
||||
return 1;
|
||||
}
|
||||
|
||||
int verify_header_twofish (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2)
|
||||
int verify_header_twofish (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2)
|
||||
{
|
||||
u32 sk_twofish[4];
|
||||
u32 lk_twofish[40];
|
||||
@ -321,7 +321,7 @@ int verify_header_twofish (__global tc_t *esalt_bufs, const u32 *ukey1, const u3
|
||||
|
||||
// 1024 bit
|
||||
|
||||
int verify_header_aes_twofish (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
int verify_header_aes_twofish (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
{
|
||||
u32 ks_aes[60];
|
||||
|
||||
@ -384,7 +384,7 @@ int verify_header_aes_twofish (__global tc_t *esalt_bufs, const u32 *ukey1, cons
|
||||
return 1;
|
||||
}
|
||||
|
||||
int verify_header_serpent_aes (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
int verify_header_serpent_aes (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
{
|
||||
u32 ks_serpent[140];
|
||||
u32 ks_aes[60];
|
||||
@ -445,7 +445,7 @@ int verify_header_serpent_aes (__global tc_t *esalt_bufs, const u32 *ukey1, cons
|
||||
return 1;
|
||||
}
|
||||
|
||||
int verify_header_twofish_serpent (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4)
|
||||
int verify_header_twofish_serpent (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4)
|
||||
{
|
||||
u32 sk_twofish[4];
|
||||
u32 lk_twofish[40];
|
||||
@ -510,7 +510,7 @@ int verify_header_twofish_serpent (__global tc_t *esalt_bufs, const u32 *ukey1,
|
||||
|
||||
// 1536 bit
|
||||
|
||||
int verify_header_aes_twofish_serpent (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, const u32 *ukey5, const u32 *ukey6, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
int verify_header_aes_twofish_serpent (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, const u32 *ukey5, const u32 *ukey6, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
{
|
||||
u32 ks_aes[60];
|
||||
|
||||
@ -579,7 +579,7 @@ int verify_header_aes_twofish_serpent (__global tc_t *esalt_bufs, const u32 *uke
|
||||
return 1;
|
||||
}
|
||||
|
||||
int verify_header_serpent_twofish_aes (__global tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, const u32 *ukey5, const u32 *ukey6, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
int verify_header_serpent_twofish_aes (__global const tc_t *esalt_bufs, const u32 *ukey1, const u32 *ukey2, const u32 *ukey3, const u32 *ukey4, const u32 *ukey5, const u32 *ukey6, SHM_TYPE u32 *s_te0, SHM_TYPE u32 *s_te1, SHM_TYPE u32 *s_te2, SHM_TYPE u32 *s_te3, SHM_TYPE u32 *s_te4, SHM_TYPE u32 *s_td0, SHM_TYPE u32 *s_td1, SHM_TYPE u32 *s_td2, SHM_TYPE u32 *s_td3, SHM_TYPE u32 *s_td4)
|
||||
{
|
||||
u32 ks_serpent[140];
|
||||
|
||||
|
||||
@ -33,14 +33,14 @@ typedef VTYPE(uint, VECT_SIZE) u32x;
|
||||
typedef VTYPE(ulong, VECT_SIZE) u64x;
|
||||
#endif
|
||||
|
||||
inline u32 l32_from_64_S (u64 a)
|
||||
u32 l32_from_64_S (u64 a)
|
||||
{
|
||||
const u32 r = (u32) (a);
|
||||
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u32 h32_from_64_S (u64 a)
|
||||
u32 h32_from_64_S (u64 a)
|
||||
{
|
||||
a >>= 32;
|
||||
|
||||
@ -49,12 +49,12 @@ inline u32 h32_from_64_S (u64 a)
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u64 hl32_to_64_S (const u32 a, const u32 b)
|
||||
u64 hl32_to_64_S (const u32 a, const u32 b)
|
||||
{
|
||||
return as_ulong ((uint2) (b, a));
|
||||
}
|
||||
|
||||
inline u32x l32_from_64 (u64x a)
|
||||
u32x l32_from_64 (u64x a)
|
||||
{
|
||||
u32x r;
|
||||
|
||||
@ -93,7 +93,7 @@ inline u32x l32_from_64 (u64x a)
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u32x h32_from_64 (u64x a)
|
||||
u32x h32_from_64 (u64x a)
|
||||
{
|
||||
a >>= 32;
|
||||
|
||||
@ -134,7 +134,7 @@ inline u32x h32_from_64 (u64x a)
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u64x hl32_to_64 (const u32x a, const u32x b)
|
||||
u64x hl32_to_64 (const u32x a, const u32x b)
|
||||
{
|
||||
u64x r;
|
||||
|
||||
@ -174,140 +174,122 @@ inline u64x hl32_to_64 (const u32x a, const u32x b)
|
||||
}
|
||||
|
||||
#ifdef IS_AMD
|
||||
inline u32 swap32_S (const u32 v)
|
||||
u32 swap32_S (const u32 v)
|
||||
{
|
||||
return (as_uint (as_uchar4 (v).s3210));
|
||||
return bitselect (rotate (v, 24u), rotate (v, 8u), 0x00ff00ffu);
|
||||
}
|
||||
|
||||
inline u64 swap64_S (const u64 v)
|
||||
u64 swap64_S (const u64 v)
|
||||
{
|
||||
return (as_ulong (as_uchar8 (v).s76543210));
|
||||
return bitselect (bitselect (rotate (v, 24ul),
|
||||
rotate (v, 8ul), 0x000000ff000000fful),
|
||||
bitselect (rotate (v, 56ul),
|
||||
rotate (v, 40ul), 0x00ff000000ff0000ul),
|
||||
0xffff0000ffff0000ul);
|
||||
}
|
||||
|
||||
inline u32 rotr32_S (const u32 a, const u32 n)
|
||||
u32 rotr32_S (const u32 a, const u32 n)
|
||||
{
|
||||
return rotate (a, 32 - n);
|
||||
return rotate (a, (32 - n));
|
||||
}
|
||||
|
||||
inline u32 rotl32_S (const u32 a, const u32 n)
|
||||
u32 rotl32_S (const u32 a, const u32 n)
|
||||
{
|
||||
return rotate (a, n);
|
||||
}
|
||||
|
||||
inline u64 rotr64_S (const u64 a, const u32 n)
|
||||
u64 rotr64_S (const u64 a, const u32 n)
|
||||
{
|
||||
const u32 a0 = h32_from_64_S (a);
|
||||
const u32 a1 = l32_from_64_S (a);
|
||||
|
||||
const u32 t0 = (n >= 32) ? amd_bitalign (a0, a1, n - 32) : amd_bitalign (a1, a0, n);
|
||||
const u32 t1 = (n >= 32) ? amd_bitalign (a1, a0, n - 32) : amd_bitalign (a0, a1, n);
|
||||
|
||||
const u64 r = hl32_to_64_S (t0, t1);
|
||||
|
||||
return r;
|
||||
return rotate (a, (u64) (64 - n));
|
||||
}
|
||||
|
||||
inline u64 rotl64_S (const u64 a, const u32 n)
|
||||
u64 rotl64_S (const u64 a, const u32 n)
|
||||
{
|
||||
return rotr64_S (a, 64 - n);
|
||||
return rotate (a, (u64) n);
|
||||
}
|
||||
|
||||
inline u32x swap32 (const u32x v)
|
||||
u32x swap32 (const u32x v)
|
||||
{
|
||||
return ((v >> 24) & 0x000000ff)
|
||||
| ((v >> 8) & 0x0000ff00)
|
||||
| ((v << 8) & 0x00ff0000)
|
||||
| ((v << 24) & 0xff000000);
|
||||
return bitselect (rotate (v, 24u), rotate (v, 8u), 0x00ff00ffu);
|
||||
}
|
||||
|
||||
inline u64x swap64 (const u64x v)
|
||||
u64x swap64 (const u64x v)
|
||||
{
|
||||
return ((v >> 56) & 0x00000000000000ff)
|
||||
| ((v >> 40) & 0x000000000000ff00)
|
||||
| ((v >> 24) & 0x0000000000ff0000)
|
||||
| ((v >> 8) & 0x00000000ff000000)
|
||||
| ((v << 8) & 0x000000ff00000000)
|
||||
| ((v << 24) & 0x0000ff0000000000)
|
||||
| ((v << 40) & 0x00ff000000000000)
|
||||
| ((v << 56) & 0xff00000000000000);
|
||||
return bitselect (bitselect (rotate (v, 24ul),
|
||||
rotate (v, 8ul), 0x000000ff000000fful),
|
||||
bitselect (rotate (v, 56ul),
|
||||
rotate (v, 40ul), 0x00ff000000ff0000ul),
|
||||
0xffff0000ffff0000ul);
|
||||
}
|
||||
|
||||
inline u32x rotr32 (const u32x a, const u32 n)
|
||||
u32x rotr32 (const u32x a, const u32 n)
|
||||
{
|
||||
return rotate (a, 32 - n);
|
||||
return rotate (a, (32 - n));
|
||||
}
|
||||
|
||||
inline u32x rotl32 (const u32x a, const u32 n)
|
||||
u32x rotl32 (const u32x a, const u32 n)
|
||||
{
|
||||
return rotate (a, n);
|
||||
}
|
||||
|
||||
inline u64x rotr64 (const u64x a, const u32 n)
|
||||
u64x rotr64 (const u64x a, const u32 n)
|
||||
{
|
||||
const u32x a0 = h32_from_64 (a);
|
||||
const u32x a1 = l32_from_64 (a);
|
||||
|
||||
const u32x t0 = (n >= 32) ? amd_bitalign (a0, a1, n - 32) : amd_bitalign (a1, a0, n);
|
||||
const u32x t1 = (n >= 32) ? amd_bitalign (a1, a0, n - 32) : amd_bitalign (a0, a1, n);
|
||||
|
||||
const u64x r = hl32_to_64 (t0, t1);
|
||||
|
||||
return r;
|
||||
return rotate (a, (u64x) (64 - n));
|
||||
}
|
||||
|
||||
inline u64x rotl64 (const u64x a, const u32 n)
|
||||
u64x rotl64 (const u64x a, const u32 n)
|
||||
{
|
||||
return rotr64 (a, 64 - n);
|
||||
return rotate (a, (u64x) n);
|
||||
}
|
||||
|
||||
inline u32x __bfe (const u32x a, const u32x b, const u32x c)
|
||||
u32x __bfe (const u32x a, const u32x b, const u32x c)
|
||||
{
|
||||
return amd_bfe (a, b, c);
|
||||
}
|
||||
|
||||
inline u32 __bfe_S (const u32 a, const u32 b, const u32 c)
|
||||
u32 __bfe_S (const u32 a, const u32 b, const u32 c)
|
||||
{
|
||||
return amd_bfe (a, b, c);
|
||||
}
|
||||
|
||||
inline u32 amd_bytealign_S (const u32 a, const u32 b, const u32 c)
|
||||
u32 amd_bytealign_S (const u32 a, const u32 b, const u32 c)
|
||||
{
|
||||
return amd_bytealign (a, b, c);
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef IS_NV
|
||||
inline u32 swap32_S (const u32 v)
|
||||
u32 swap32_S (const u32 v)
|
||||
{
|
||||
return (as_uint (as_uchar4 (v).s3210));
|
||||
}
|
||||
|
||||
inline u64 swap64_S (const u64 v)
|
||||
u64 swap64_S (const u64 v)
|
||||
{
|
||||
return (as_ulong (as_uchar8 (v).s76543210));
|
||||
}
|
||||
|
||||
inline u32 rotr32_S (const u32 a, const u32 n)
|
||||
u32 rotr32_S (const u32 a, const u32 n)
|
||||
{
|
||||
return rotate (a, 32 - n);
|
||||
return rotate (a, (32 - n));
|
||||
}
|
||||
|
||||
inline u32 rotl32_S (const u32 a, const u32 n)
|
||||
u32 rotl32_S (const u32 a, const u32 n)
|
||||
{
|
||||
return rotate (a, n);
|
||||
}
|
||||
|
||||
inline u64 rotr64_S (const u64 a, const u32 n)
|
||||
u64 rotr64_S (const u64 a, const u32 n)
|
||||
{
|
||||
return rotate (a, (u64) 64 - n);
|
||||
return rotate (a, (u64) (64 - n));
|
||||
}
|
||||
|
||||
inline u64 rotl64_S (const u64 a, const u32 n)
|
||||
u64 rotl64_S (const u64 a, const u32 n)
|
||||
{
|
||||
return rotr64_S (a, 64 - n);
|
||||
return rotate (a, (u64) n);
|
||||
}
|
||||
|
||||
inline u32x swap32 (const u32x v)
|
||||
u32x swap32 (const u32x v)
|
||||
{
|
||||
return ((v >> 24) & 0x000000ff)
|
||||
| ((v >> 8) & 0x0000ff00)
|
||||
@ -315,7 +297,7 @@ inline u32x swap32 (const u32x v)
|
||||
| ((v << 24) & 0xff000000);
|
||||
}
|
||||
|
||||
inline u64x swap64 (const u64x v)
|
||||
u64x swap64 (const u64x v)
|
||||
{
|
||||
return ((v >> 56) & 0x00000000000000ff)
|
||||
| ((v >> 40) & 0x000000000000ff00)
|
||||
@ -327,27 +309,27 @@ inline u64x swap64 (const u64x v)
|
||||
| ((v << 56) & 0xff00000000000000);
|
||||
}
|
||||
|
||||
inline u32x rotr32 (const u32x a, const u32 n)
|
||||
u32x rotr32 (const u32x a, const u32 n)
|
||||
{
|
||||
return rotate (a, 32 - n);
|
||||
return rotate (a, (32 - n));
|
||||
}
|
||||
|
||||
inline u32x rotl32 (const u32x a, const u32 n)
|
||||
u32x rotl32 (const u32x a, const u32 n)
|
||||
{
|
||||
return rotate (a, n);
|
||||
}
|
||||
|
||||
inline u64x rotr64 (const u64x a, const u32 n)
|
||||
u64x rotr64 (const u64x a, const u32 n)
|
||||
{
|
||||
return rotate (a, (u64) 64 - n);
|
||||
return rotate (a, (u64x) (64 - n));
|
||||
}
|
||||
|
||||
inline u64x rotl64 (const u64x a, const u32 n)
|
||||
u64x rotl64 (const u64x a, const u32 n)
|
||||
{
|
||||
return rotate (a, (u64) n);
|
||||
return rotate (a, (u64x) n);
|
||||
}
|
||||
|
||||
inline u32x __byte_perm (const u32x a, const u32x b, const u32x c)
|
||||
u32x __byte_perm (const u32x a, const u32x b, const u32x c)
|
||||
{
|
||||
u32x r;
|
||||
|
||||
@ -386,7 +368,7 @@ inline u32x __byte_perm (const u32x a, const u32x b, const u32x c)
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u32 __byte_perm_S (const u32 a, const u32 b, const u32 c)
|
||||
u32 __byte_perm_S (const u32 a, const u32 b, const u32 c)
|
||||
{
|
||||
u32 r;
|
||||
|
||||
@ -395,7 +377,7 @@ inline u32 __byte_perm_S (const u32 a, const u32 b, const u32 c)
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u32x __bfe (const u32x a, const u32x b, const u32x c)
|
||||
u32x __bfe (const u32x a, const u32x b, const u32x c)
|
||||
{
|
||||
u32x r;
|
||||
|
||||
@ -434,7 +416,7 @@ inline u32x __bfe (const u32x a, const u32x b, const u32x c)
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u32 __bfe_S (const u32 a, const u32 b, const u32 c)
|
||||
u32 __bfe_S (const u32 a, const u32 b, const u32 c)
|
||||
{
|
||||
u32 r;
|
||||
|
||||
@ -443,7 +425,7 @@ inline u32 __bfe_S (const u32 a, const u32 b, const u32 c)
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u32x amd_bytealign (const u32x a, const u32x b, const u32x c)
|
||||
u32x amd_bytealign (const u32x a, const u32x b, const u32x c)
|
||||
{
|
||||
u32x r;
|
||||
|
||||
@ -490,7 +472,7 @@ inline u32x amd_bytealign (const u32x a, const u32x b, const u32x c)
|
||||
return r;
|
||||
}
|
||||
|
||||
inline u32 amd_bytealign_S (const u32 a, const u32 b, const u32 c)
|
||||
u32 amd_bytealign_S (const u32 a, const u32 b, const u32 c)
|
||||
{
|
||||
u32 r;
|
||||
|
||||
@ -509,37 +491,37 @@ inline u32 amd_bytealign_S (const u32 a, const u32 b, const u32 c)
|
||||
#endif
|
||||
|
||||
#ifdef IS_GENERIC
|
||||
inline u32 swap32_S (const u32 v)
|
||||
u32 swap32_S (const u32 v)
|
||||
{
|
||||
return (as_uint (as_uchar4 (v).s3210));
|
||||
}
|
||||
|
||||
inline u64 swap64_S (const u64 v)
|
||||
u64 swap64_S (const u64 v)
|
||||
{
|
||||
return (as_ulong (as_uchar8 (v).s76543210));
|
||||
}
|
||||
|
||||
inline u32 rotr32_S (const u32 a, const u32 n)
|
||||
u32 rotr32_S (const u32 a, const u32 n)
|
||||
{
|
||||
return rotate (a, 32 - n);
|
||||
return rotate (a, (32 - n));
|
||||
}
|
||||
|
||||
inline u32 rotl32_S (const u32 a, const u32 n)
|
||||
u32 rotl32_S (const u32 a, const u32 n)
|
||||
{
|
||||
return rotate (a, n);
|
||||
}
|
||||
|
||||
inline u64 rotr64_S (const u64 a, const u32 n)
|
||||
u64 rotr64_S (const u64 a, const u32 n)
|
||||
{
|
||||
return rotate (a, (u64) 64 - n);
|
||||
return rotate (a, (u64) (64 - n));
|
||||
}
|
||||
|
||||
inline u64 rotl64_S (const u64 a, const u32 n)
|
||||
u64 rotl64_S (const u64 a, const u32 n)
|
||||
{
|
||||
return rotate (a, (u64) n);
|
||||
}
|
||||
|
||||
inline u32x swap32 (const u32x v)
|
||||
u32x swap32 (const u32x v)
|
||||
{
|
||||
return ((v >> 24) & 0x000000ff)
|
||||
| ((v >> 8) & 0x0000ff00)
|
||||
@ -547,7 +529,7 @@ inline u32x swap32 (const u32x v)
|
||||
| ((v << 24) & 0xff000000);
|
||||
}
|
||||
|
||||
inline u64x swap64 (const u64x v)
|
||||
u64x swap64 (const u64x v)
|
||||
{
|
||||
return ((v >> 56) & 0x00000000000000ff)
|
||||
| ((v >> 40) & 0x000000000000ff00)
|
||||
@ -559,27 +541,27 @@ inline u64x swap64 (const u64x v)
|
||||
| ((v << 56) & 0xff00000000000000);
|
||||
}
|
||||
|
||||
inline u32x rotr32 (const u32x a, const u32 n)
|
||||
u32x rotr32 (const u32x a, const u32 n)
|
||||
{
|
||||
return rotate (a, 32 - n);
|
||||
return rotate (a, (32 - n));
|
||||
}
|
||||
|
||||
inline u32x rotl32 (const u32x a, const u32 n)
|
||||
u32x rotl32 (const u32x a, const u32 n)
|
||||
{
|
||||
return rotate (a, n);
|
||||
}
|
||||
|
||||
inline u64x rotr64 (const u64x a, const u32 n)
|
||||
u64x rotr64 (const u64x a, const u32 n)
|
||||
{
|
||||
return rotate (a, (u64) 64 - n);
|
||||
return rotate (a, (u64x) (64 - n));
|
||||
}
|
||||
|
||||
inline u64x rotl64 (const u64x a, const u32 n)
|
||||
u64x rotl64 (const u64x a, const u32 n)
|
||||
{
|
||||
return rotate (a, (u64) n);
|
||||
return rotate (a, (u64x) n);
|
||||
}
|
||||
|
||||
inline u32x __bfe (const u32x a, const u32x b, const u32x c)
|
||||
u32x __bfe (const u32x a, const u32x b, const u32x c)
|
||||
{
|
||||
#define BIT(x) ((u32x) (1u) << (x))
|
||||
#define BIT_MASK(x) (BIT (x) - 1)
|
||||
@ -592,7 +574,7 @@ inline u32x __bfe (const u32x a, const u32x b, const u32x c)
|
||||
#undef BFE
|
||||
}
|
||||
|
||||
inline u32 __bfe_S (const u32 a, const u32 b, const u32 c)
|
||||
u32 __bfe_S (const u32 a, const u32 b, const u32 c)
|
||||
{
|
||||
#define BIT(x) (1u << (x))
|
||||
#define BIT_MASK(x) (BIT (x) - 1)
|
||||
@ -605,7 +587,7 @@ inline u32 __bfe_S (const u32 a, const u32 b, const u32 c)
|
||||
#undef BFE
|
||||
}
|
||||
|
||||
inline u32x amd_bytealign (const u32x a, const u32x b, const u32 c)
|
||||
u32x amd_bytealign (const u32x a, const u32x b, const u32 c)
|
||||
{
|
||||
#if VECT_SIZE == 1
|
||||
const u64x tmp = ((((u64x) (a)) << 32) | ((u64x) (b))) >> ((c & 3) * 8);
|
||||
@ -638,7 +620,7 @@ inline u32x amd_bytealign (const u32x a, const u32x b, const u32 c)
|
||||
#endif
|
||||
}
|
||||
|
||||
inline u32 amd_bytealign_S (const u32 a, const u32 b, const u32 c)
|
||||
u32 amd_bytealign_S (const u32 a, const u32 b, const u32 c)
|
||||
{
|
||||
const u64 tmp = ((((u64) a) << 32) | ((u64) b)) >> ((c & 3) * 8);
|
||||
|
||||
@ -809,10 +791,10 @@ typedef struct bitcoin_wallet
|
||||
|
||||
typedef struct sip
|
||||
{
|
||||
u32 salt_buf[30];
|
||||
u32 salt_buf[32];
|
||||
u32 salt_len;
|
||||
|
||||
u32 esalt_buf[38];
|
||||
u32 esalt_buf[48];
|
||||
u32 esalt_len;
|
||||
|
||||
} sip_t;
|
||||
|
||||
@ -153,9 +153,6 @@
|
||||
#if KERN_TYPE == 13800
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 14100
|
||||
#undef _unroll
|
||||
#endif
|
||||
|
||||
// nvidia specific
|
||||
|
||||
@ -177,6 +174,9 @@
|
||||
#if KERN_TYPE == 14000
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 14100
|
||||
#undef _unroll
|
||||
#endif
|
||||
|
||||
#endif
|
||||
#endif
|
||||
@ -186,25 +186,7 @@
|
||||
#ifdef IS_AMD
|
||||
#ifdef IS_GPU
|
||||
|
||||
#if KERN_TYPE == 1700
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 1710
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 5200
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 10800
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 10900
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 12800
|
||||
#undef _unroll
|
||||
#endif
|
||||
#if KERN_TYPE == 12900
|
||||
#if KERN_TYPE == 8000
|
||||
#undef _unroll
|
||||
#endif
|
||||
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -104,8 +102,6 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -108,8 +106,6 @@ __kernel void m00000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -119,8 +115,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
@ -106,8 +104,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00010_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -123,8 +119,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -136,8 +130,6 @@ __kernel void m00010_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
@ -108,8 +106,6 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00020_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
@ -114,8 +112,6 @@ __kernel void m00020_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -119,8 +115,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
@ -106,8 +104,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00030_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -123,8 +119,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -136,8 +130,6 @@ __kernel void m00030_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
@ -108,8 +106,6 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00040_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
@ -114,8 +112,6 @@ __kernel void m00040_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_ctx_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -117,8 +113,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -130,8 +124,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -82,13 +78,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];
|
||||
@ -135,8 +131,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -148,8 +142,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -180,13 +172,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];
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00050_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -121,8 +117,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -134,8 +128,6 @@ __kernel void m00050_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_hmac_ctx_t ctx0;
|
||||
@ -119,8 +115,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_hmac_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_hmac_ctx_t ctx0;
|
||||
@ -137,8 +133,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -150,8 +144,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_hmac_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00060_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_hmac_ctx_vector_t ctx0;
|
||||
@ -123,8 +119,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -136,8 +130,6 @@ __kernel void m00060_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md5_hmac_ctx_vector_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -104,8 +102,6 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -108,8 +106,6 @@ __kernel void m00100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -119,8 +115,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
@ -106,8 +104,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00110_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -123,8 +119,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -136,8 +130,6 @@ __kernel void m00110_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
@ -108,8 +106,6 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00120_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
@ -114,8 +112,6 @@ __kernel void m00120_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -119,8 +115,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
@ -106,8 +104,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00130_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -124,8 +120,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -137,8 +131,6 @@ __kernel void m00130_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
@ -108,8 +106,6 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00140_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
@ -114,8 +112,6 @@ __kernel void m00140_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_ctx_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -117,8 +113,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -130,8 +124,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -82,13 +78,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];
|
||||
@ -135,8 +131,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -148,8 +142,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -180,13 +172,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];
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00150_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -121,8 +117,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -134,8 +128,6 @@ __kernel void m00150_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_hmac_ctx_t ctx0;
|
||||
@ -119,8 +115,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_hmac_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_hmac_ctx_t ctx0;
|
||||
@ -137,8 +133,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -150,8 +144,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_hmac_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m00160_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_hmac_ctx_vector_t ctx0;
|
||||
@ -123,8 +119,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -136,8 +130,6 @@ __kernel void m00160_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha1_hmac_ctx_vector_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -127,8 +125,6 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -131,8 +129,6 @@ __kernel void m00300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -76,8 +76,6 @@ __kernel void m00400_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
u32 digest[4];
|
||||
|
||||
@ -10,111 +10,14 @@
|
||||
#include "inc_hash_functions.cl"
|
||||
#include "inc_types.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
|
||||
#define md5crypt_magic 0x00243124u
|
||||
|
||||
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 = 0;
|
||||
|
||||
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 memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[4])
|
||||
void memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 offset, const u32 append[4])
|
||||
{
|
||||
u32 tmp0;
|
||||
u32 tmp1;
|
||||
@ -122,44 +25,45 @@ void memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const
|
||||
u32 tmp3;
|
||||
u32 tmp4;
|
||||
|
||||
const int offset_mod_4 = offset & 3;
|
||||
|
||||
const int offset_minus_4 = 4 - offset_mod_4;
|
||||
|
||||
#if defined IS_AMD || defined IS_GENERIC
|
||||
u32 in0 = swap32_S (append[0]);
|
||||
u32 in1 = swap32_S (append[1]);
|
||||
u32 in2 = swap32_S (append[2]);
|
||||
u32 in3 = swap32_S (append[3]);
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
|
||||
tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
|
||||
tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
|
||||
tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
|
||||
tmp4 = amd_bytealign ( 0, append[3], offset_minus_4);
|
||||
|
||||
const u32 mod = block_len & 3;
|
||||
|
||||
if (mod == 0)
|
||||
{
|
||||
tmp0 = tmp1;
|
||||
tmp1 = tmp2;
|
||||
tmp2 = tmp3;
|
||||
tmp3 = tmp4;
|
||||
tmp4 = 0;
|
||||
}
|
||||
tmp0 = amd_bytealign ( 0, in0, offset);
|
||||
tmp1 = amd_bytealign (in0, in1, offset);
|
||||
tmp2 = amd_bytealign (in1, in2, offset);
|
||||
tmp3 = amd_bytealign (in2, in3, offset);
|
||||
tmp4 = amd_bytealign (in3, 0, offset);
|
||||
|
||||
tmp0 = swap32_S (tmp0);
|
||||
tmp1 = swap32_S (tmp1);
|
||||
tmp2 = swap32_S (tmp2);
|
||||
tmp3 = swap32_S (tmp3);
|
||||
tmp4 = swap32_S (tmp4);
|
||||
#endif
|
||||
|
||||
#ifdef IS_NV
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
|
||||
|
||||
tmp0 = __byte_perm ( 0, append[0], selector);
|
||||
tmp1 = __byte_perm (append[0], append[1], selector);
|
||||
tmp2 = __byte_perm (append[1], append[2], selector);
|
||||
tmp3 = __byte_perm (append[2], append[3], selector);
|
||||
tmp4 = __byte_perm (append[3], 0, selector);
|
||||
u32 in0 = append[0];
|
||||
u32 in1 = append[1];
|
||||
u32 in2 = append[2];
|
||||
u32 in3 = append[3];
|
||||
|
||||
tmp0 = __byte_perm ( 0, in0, selector);
|
||||
tmp1 = __byte_perm (in0, in1, selector);
|
||||
tmp2 = __byte_perm (in1, in2, selector);
|
||||
tmp3 = __byte_perm (in2, in3, selector);
|
||||
tmp4 = __byte_perm (in3, 0, selector);
|
||||
#endif
|
||||
|
||||
const u32 div = block_len / 4;
|
||||
const u32 div = offset / 4;
|
||||
|
||||
switch (div)
|
||||
{
|
||||
@ -226,7 +130,7 @@ void memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const
|
||||
}
|
||||
}
|
||||
|
||||
void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[4])
|
||||
void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 offset, const u32 append[4])
|
||||
{
|
||||
u32 tmp0;
|
||||
u32 tmp1;
|
||||
@ -234,44 +138,47 @@ void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], c
|
||||
u32 tmp3;
|
||||
u32 tmp4;
|
||||
|
||||
const int offset_mod_4 = offset & 3;
|
||||
|
||||
const int offset_minus_4 = 4 - offset_mod_4;
|
||||
|
||||
#if defined IS_AMD || defined IS_GENERIC
|
||||
u32 in0 = swap32_S (append[0]);
|
||||
u32 in1 = swap32_S (append[1]);
|
||||
u32 in2 = swap32_S (append[2]);
|
||||
u32 in3 = swap32_S (append[3]);
|
||||
u32 in4 = 0x80000000;
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
|
||||
tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
|
||||
tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
|
||||
tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
|
||||
tmp4 = amd_bytealign ( 0x80, append[3], offset_minus_4);
|
||||
|
||||
const u32 mod = block_len & 3;
|
||||
|
||||
if (mod == 0)
|
||||
{
|
||||
tmp0 = tmp1;
|
||||
tmp1 = tmp2;
|
||||
tmp2 = tmp3;
|
||||
tmp3 = tmp4;
|
||||
tmp4 = 0x80;
|
||||
}
|
||||
tmp0 = amd_bytealign ( 0, in0, offset);
|
||||
tmp1 = amd_bytealign (in0, in1, offset);
|
||||
tmp2 = amd_bytealign (in1, in2, offset);
|
||||
tmp3 = amd_bytealign (in2, in3, offset);
|
||||
tmp4 = amd_bytealign (in3, in4, offset);
|
||||
|
||||
tmp0 = swap32_S (tmp0);
|
||||
tmp1 = swap32_S (tmp1);
|
||||
tmp2 = swap32_S (tmp2);
|
||||
tmp3 = swap32_S (tmp3);
|
||||
tmp4 = swap32_S (tmp4);
|
||||
#endif
|
||||
|
||||
#ifdef IS_NV
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
|
||||
|
||||
tmp0 = __byte_perm ( 0, append[0], selector);
|
||||
tmp1 = __byte_perm (append[0], append[1], selector);
|
||||
tmp2 = __byte_perm (append[1], append[2], selector);
|
||||
tmp3 = __byte_perm (append[2], append[3], selector);
|
||||
tmp4 = __byte_perm (append[3], 0x80, selector);
|
||||
u32 in0 = append[0];
|
||||
u32 in1 = append[1];
|
||||
u32 in2 = append[2];
|
||||
u32 in3 = append[3];
|
||||
u32 in4 = 0x80;
|
||||
|
||||
tmp0 = __byte_perm ( 0, in0, selector);
|
||||
tmp1 = __byte_perm (in0, in1, selector);
|
||||
tmp2 = __byte_perm (in1, in2, selector);
|
||||
tmp3 = __byte_perm (in2, in3, selector);
|
||||
tmp4 = __byte_perm (in3, in4, selector);
|
||||
#endif
|
||||
|
||||
const u32 div = block_len / 4;
|
||||
const u32 div = offset / 4;
|
||||
|
||||
switch (div)
|
||||
{
|
||||
@ -338,44 +245,41 @@ void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], c
|
||||
}
|
||||
}
|
||||
|
||||
void memcat8 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[2])
|
||||
void memcat8 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 offset, const u32 append[2])
|
||||
{
|
||||
u32 tmp0;
|
||||
u32 tmp1;
|
||||
u32 tmp2;
|
||||
|
||||
const int offset_mod_4 = offset & 3;
|
||||
|
||||
const int offset_minus_4 = 4 - offset_mod_4;
|
||||
|
||||
#if defined IS_AMD || defined IS_GENERIC
|
||||
u32 in0 = swap32_S (append[0]);
|
||||
u32 in1 = swap32_S (append[1]);
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
|
||||
tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
|
||||
tmp2 = amd_bytealign ( 0, append[1], offset_minus_4);
|
||||
|
||||
const u32 mod = block_len & 3;
|
||||
|
||||
if (mod == 0)
|
||||
{
|
||||
tmp0 = tmp1;
|
||||
tmp1 = tmp2;
|
||||
tmp2 = 0;
|
||||
}
|
||||
tmp0 = amd_bytealign ( 0, in0, offset);
|
||||
tmp1 = amd_bytealign (in0, in1, offset);
|
||||
tmp2 = amd_bytealign (in1, 0, offset);
|
||||
|
||||
tmp0 = swap32_S (tmp0);
|
||||
tmp1 = swap32_S (tmp1);
|
||||
tmp2 = swap32_S (tmp2);
|
||||
#endif
|
||||
|
||||
#ifdef IS_NV
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
|
||||
|
||||
tmp0 = __byte_perm ( 0, append[0], selector);
|
||||
tmp1 = __byte_perm (append[0], append[1], selector);
|
||||
tmp2 = __byte_perm (append[1], 0, selector);
|
||||
u32 in0 = append[0];
|
||||
u32 in1 = append[1];
|
||||
|
||||
tmp0 = __byte_perm ( 0, in0, selector);
|
||||
tmp1 = __byte_perm (in0, in1, selector);
|
||||
tmp2 = __byte_perm (in1, 0, selector);
|
||||
#endif
|
||||
|
||||
const u32 div = block_len / 4;
|
||||
const u32 div = offset / 4;
|
||||
|
||||
switch (div)
|
||||
{
|
||||
|
||||
@ -40,8 +40,6 @@ __kernel void m00500_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -53,8 +51,6 @@ __kernel void m00500_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -154,8 +150,6 @@ __kernel void m00500_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -167,8 +161,6 @@ __kernel void m00500_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -104,8 +102,6 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m00900_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -108,8 +106,6 @@ __kernel void m00900_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -104,8 +102,6 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01000_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -108,8 +106,6 @@ __kernel void m01000_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,21 +39,17 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
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++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -132,21 +128,17 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
|
||||
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++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -32,13 +32,11 @@ __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++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md4_ctx_t ctx0;
|
||||
@ -114,13 +112,11 @@ __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++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
md4_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01100_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -136,8 +132,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -149,8 +143,6 @@ __kernel void m01100_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -104,8 +102,6 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01300_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -108,8 +106,6 @@ __kernel void m01300_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -104,8 +102,6 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01400_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -108,8 +106,6 @@ __kernel void m01400_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -119,8 +115,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
@ -106,8 +104,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01410_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -123,8 +119,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -136,8 +130,6 @@ __kernel void m01410_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
@ -108,8 +106,6 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01420_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
@ -114,8 +112,6 @@ __kernel void m01420_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -119,8 +115,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
@ -106,8 +104,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01430_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -124,8 +120,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -137,8 +131,6 @@ __kernel void m01430_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
@ -108,8 +106,6 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01440_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
@ -114,8 +112,6 @@ __kernel void m01440_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_ctx_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -117,8 +113,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -130,8 +124,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -82,13 +78,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];
|
||||
@ -135,8 +131,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -148,8 +142,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -180,13 +172,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];
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01450_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -121,8 +117,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -134,8 +128,6 @@ __kernel void m01450_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_hmac_ctx_t ctx0;
|
||||
@ -119,8 +115,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_hmac_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_hmac_ctx_t ctx0;
|
||||
@ -137,8 +133,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = swap32_S (pws[gid].i[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -150,8 +144,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_hmac_ctx_t ctx0;
|
||||
|
||||
@ -221,7 +221,7 @@ void hmac_sha256_run (u32x w0[4], u32x w1[4], u32x w2[4], u32x w3[4], u32x ipad[
|
||||
sha256_transform (w0, w1, w2, w3, digest);
|
||||
}
|
||||
|
||||
void m01460m (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 void *esal_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)
|
||||
void m01460m (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)
|
||||
{
|
||||
/**
|
||||
* modifier
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01460_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_hmac_ctx_vector_t ctx0;
|
||||
@ -123,8 +119,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -136,8 +130,6 @@ __kernel void m01460_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha256_hmac_ctx_vector_t ctx0;
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@ -8,6 +8,7 @@
|
||||
#include "inc_hash_functions.cl"
|
||||
#include "inc_types.cl"
|
||||
#include "inc_common.cl"
|
||||
#include "inc_hash_md5.cl"
|
||||
|
||||
#define COMPARE_S "inc_comp_single.cl"
|
||||
#define COMPARE_M "inc_comp_multi.cl"
|
||||
@ -15,105 +16,7 @@
|
||||
#define md5apr1_magic0 0x72706124u
|
||||
#define md5apr1_magic1 0x00002431u
|
||||
|
||||
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 = 0;
|
||||
|
||||
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 memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[4])
|
||||
void memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 offset, const u32 append[4])
|
||||
{
|
||||
u32 tmp0;
|
||||
u32 tmp1;
|
||||
@ -121,44 +24,45 @@ void memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const
|
||||
u32 tmp3;
|
||||
u32 tmp4;
|
||||
|
||||
const int offset_mod_4 = offset & 3;
|
||||
|
||||
const int offset_minus_4 = 4 - offset_mod_4;
|
||||
|
||||
#if defined IS_AMD || defined IS_GENERIC
|
||||
u32 in0 = swap32_S (append[0]);
|
||||
u32 in1 = swap32_S (append[1]);
|
||||
u32 in2 = swap32_S (append[2]);
|
||||
u32 in3 = swap32_S (append[3]);
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
|
||||
tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
|
||||
tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
|
||||
tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
|
||||
tmp4 = amd_bytealign ( 0, append[3], offset_minus_4);
|
||||
|
||||
const u32 mod = block_len & 3;
|
||||
|
||||
if (mod == 0)
|
||||
{
|
||||
tmp0 = tmp1;
|
||||
tmp1 = tmp2;
|
||||
tmp2 = tmp3;
|
||||
tmp3 = tmp4;
|
||||
tmp4 = 0;
|
||||
}
|
||||
tmp0 = amd_bytealign ( 0, in0, offset);
|
||||
tmp1 = amd_bytealign (in0, in1, offset);
|
||||
tmp2 = amd_bytealign (in1, in2, offset);
|
||||
tmp3 = amd_bytealign (in2, in3, offset);
|
||||
tmp4 = amd_bytealign (in3, 0, offset);
|
||||
|
||||
tmp0 = swap32_S (tmp0);
|
||||
tmp1 = swap32_S (tmp1);
|
||||
tmp2 = swap32_S (tmp2);
|
||||
tmp3 = swap32_S (tmp3);
|
||||
tmp4 = swap32_S (tmp4);
|
||||
#endif
|
||||
|
||||
#ifdef IS_NV
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
|
||||
|
||||
tmp0 = __byte_perm ( 0, append[0], selector);
|
||||
tmp1 = __byte_perm (append[0], append[1], selector);
|
||||
tmp2 = __byte_perm (append[1], append[2], selector);
|
||||
tmp3 = __byte_perm (append[2], append[3], selector);
|
||||
tmp4 = __byte_perm (append[3], 0, selector);
|
||||
u32 in0 = append[0];
|
||||
u32 in1 = append[1];
|
||||
u32 in2 = append[2];
|
||||
u32 in3 = append[3];
|
||||
|
||||
tmp0 = __byte_perm ( 0, in0, selector);
|
||||
tmp1 = __byte_perm (in0, in1, selector);
|
||||
tmp2 = __byte_perm (in1, in2, selector);
|
||||
tmp3 = __byte_perm (in2, in3, selector);
|
||||
tmp4 = __byte_perm (in3, 0, selector);
|
||||
#endif
|
||||
|
||||
const u32 div = block_len / 4;
|
||||
const u32 div = offset / 4;
|
||||
|
||||
switch (div)
|
||||
{
|
||||
@ -225,7 +129,7 @@ void memcat16 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const
|
||||
}
|
||||
}
|
||||
|
||||
void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[4])
|
||||
void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 offset, const u32 append[4])
|
||||
{
|
||||
u32 tmp0;
|
||||
u32 tmp1;
|
||||
@ -233,44 +137,47 @@ void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], c
|
||||
u32 tmp3;
|
||||
u32 tmp4;
|
||||
|
||||
const int offset_mod_4 = offset & 3;
|
||||
|
||||
const int offset_minus_4 = 4 - offset_mod_4;
|
||||
|
||||
#if defined IS_AMD || defined IS_GENERIC
|
||||
u32 in0 = swap32_S (append[0]);
|
||||
u32 in1 = swap32_S (append[1]);
|
||||
u32 in2 = swap32_S (append[2]);
|
||||
u32 in3 = swap32_S (append[3]);
|
||||
u32 in4 = 0x80000000;
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
|
||||
tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
|
||||
tmp2 = amd_bytealign (append[2], append[1], offset_minus_4);
|
||||
tmp3 = amd_bytealign (append[3], append[2], offset_minus_4);
|
||||
tmp4 = amd_bytealign ( 0x80, append[3], offset_minus_4);
|
||||
|
||||
const u32 mod = block_len & 3;
|
||||
|
||||
if (mod == 0)
|
||||
{
|
||||
tmp0 = tmp1;
|
||||
tmp1 = tmp2;
|
||||
tmp2 = tmp3;
|
||||
tmp3 = tmp4;
|
||||
tmp4 = 0x80;
|
||||
}
|
||||
tmp0 = amd_bytealign ( 0, in0, offset);
|
||||
tmp1 = amd_bytealign (in0, in1, offset);
|
||||
tmp2 = amd_bytealign (in1, in2, offset);
|
||||
tmp3 = amd_bytealign (in2, in3, offset);
|
||||
tmp4 = amd_bytealign (in3, in4, offset);
|
||||
|
||||
tmp0 = swap32_S (tmp0);
|
||||
tmp1 = swap32_S (tmp1);
|
||||
tmp2 = swap32_S (tmp2);
|
||||
tmp3 = swap32_S (tmp3);
|
||||
tmp4 = swap32_S (tmp4);
|
||||
#endif
|
||||
|
||||
#ifdef IS_NV
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
|
||||
|
||||
tmp0 = __byte_perm ( 0, append[0], selector);
|
||||
tmp1 = __byte_perm (append[0], append[1], selector);
|
||||
tmp2 = __byte_perm (append[1], append[2], selector);
|
||||
tmp3 = __byte_perm (append[2], append[3], selector);
|
||||
tmp4 = __byte_perm (append[3], 0x80, selector);
|
||||
u32 in0 = append[0];
|
||||
u32 in1 = append[1];
|
||||
u32 in2 = append[2];
|
||||
u32 in3 = append[3];
|
||||
u32 in4 = 0x80;
|
||||
|
||||
tmp0 = __byte_perm ( 0, in0, selector);
|
||||
tmp1 = __byte_perm (in0, in1, selector);
|
||||
tmp2 = __byte_perm (in1, in2, selector);
|
||||
tmp3 = __byte_perm (in2, in3, selector);
|
||||
tmp4 = __byte_perm (in3, in4, selector);
|
||||
#endif
|
||||
|
||||
const u32 div = block_len / 4;
|
||||
const u32 div = offset / 4;
|
||||
|
||||
switch (div)
|
||||
{
|
||||
@ -337,44 +244,41 @@ void memcat16_x80 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], c
|
||||
}
|
||||
}
|
||||
|
||||
void memcat8 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 block_len, const u32 append[2])
|
||||
void memcat8 (u32 block0[4], u32 block1[4], u32 block2[4], u32 block3[4], const u32 offset, const u32 append[2])
|
||||
{
|
||||
u32 tmp0;
|
||||
u32 tmp1;
|
||||
u32 tmp2;
|
||||
|
||||
const int offset_mod_4 = offset & 3;
|
||||
|
||||
const int offset_minus_4 = 4 - offset_mod_4;
|
||||
|
||||
#if defined IS_AMD || defined IS_GENERIC
|
||||
u32 in0 = swap32_S (append[0]);
|
||||
u32 in1 = swap32_S (append[1]);
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
tmp0 = amd_bytealign (append[0], 0, offset_minus_4);
|
||||
tmp1 = amd_bytealign (append[1], append[0], offset_minus_4);
|
||||
tmp2 = amd_bytealign ( 0, append[1], offset_minus_4);
|
||||
|
||||
const u32 mod = block_len & 3;
|
||||
|
||||
if (mod == 0)
|
||||
{
|
||||
tmp0 = tmp1;
|
||||
tmp1 = tmp2;
|
||||
tmp2 = 0;
|
||||
}
|
||||
tmp0 = amd_bytealign ( 0, in0, offset);
|
||||
tmp1 = amd_bytealign (in0, in1, offset);
|
||||
tmp2 = amd_bytealign (in1, 0, offset);
|
||||
|
||||
tmp0 = swap32_S (tmp0);
|
||||
tmp1 = swap32_S (tmp1);
|
||||
tmp2 = swap32_S (tmp2);
|
||||
#endif
|
||||
|
||||
#ifdef IS_NV
|
||||
|
||||
const int offset_minus_4 = 4 - (block_len & 3);
|
||||
|
||||
const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff;
|
||||
|
||||
tmp0 = __byte_perm ( 0, append[0], selector);
|
||||
tmp1 = __byte_perm (append[0], append[1], selector);
|
||||
tmp2 = __byte_perm (append[1], 0, selector);
|
||||
u32 in0 = append[0];
|
||||
u32 in1 = append[1];
|
||||
|
||||
tmp0 = __byte_perm ( 0, in0, selector);
|
||||
tmp1 = __byte_perm (in0, in1, selector);
|
||||
tmp2 = __byte_perm (in1, 0, selector);
|
||||
#endif
|
||||
|
||||
const u32 div = block_len / 4;
|
||||
const u32 div = offset / 4;
|
||||
|
||||
switch (div)
|
||||
{
|
||||
|
||||
@ -41,8 +41,6 @@ __kernel void m01600_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -54,8 +52,6 @@ __kernel void m01600_init (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -156,8 +152,6 @@ __kernel void m01600_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -169,8 +163,6 @@ __kernel void m01600_loop (__global pw_t *pws, __global const kernel_rule_t *rul
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = salt_bufs[salt_pos].salt_buf[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -104,8 +102,6 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01700_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -108,8 +106,6 @@ __kernel void m01700_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -119,8 +115,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
@ -106,8 +104,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01710_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -123,8 +119,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -136,8 +130,6 @@ __kernel void m01710_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
@ -108,8 +106,6 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01720_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
@ -114,8 +112,6 @@ __kernel void m01720_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -119,8 +115,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -132,8 +126,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
@ -106,8 +104,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -50,8 +48,6 @@ __kernel void m01730_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -124,8 +120,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -137,8 +131,6 @@ __kernel void m01730_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
@ -108,8 +106,6 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -37,8 +37,6 @@ __kernel void m01740_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
@ -114,8 +112,6 @@ __kernel void m01740_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
sha512_ctx_t ctx0;
|
||||
|
||||
@ -39,8 +39,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -52,8 +50,6 @@ __kernel void m01750_mxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
@ -117,8 +113,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < pw_lenv; idx++)
|
||||
{
|
||||
w[idx] = pws[gid].i[idx];
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
const u32 salt_len = salt_bufs[salt_pos].salt_len;
|
||||
@ -130,8 +124,6 @@ __kernel void m01750_sxx (__global pw_t *pws, __global const kernel_rule_t *rule
|
||||
for (int idx = 0; idx < salt_lenv; idx++)
|
||||
{
|
||||
s[idx] = swap32_S (salt_bufs[salt_pos].salt_buf[idx]);
|
||||
|
||||
barrier (CLK_GLOBAL_MEM_FENCE);
|
||||
}
|
||||
|
||||
/**
|
||||
|
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue
Block a user