|
|
|
@ -30,14 +30,7 @@ __kernel void m10800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru
|
|
|
|
|
* base
|
|
|
|
|
*/
|
|
|
|
|
|
|
|
|
|
const u32 pw_len = pws[gid].pw_len;
|
|
|
|
|
|
|
|
|
|
u32 w[64] = { 0 };
|
|
|
|
|
|
|
|
|
|
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
|
|
|
|
|
{
|
|
|
|
|
w[idx] = swap32_S (pws[gid].i[idx]);
|
|
|
|
|
}
|
|
|
|
|
COPY_PW (pws[gid]);
|
|
|
|
|
|
|
|
|
|
/**
|
|
|
|
|
* loop
|
|
|
|
@ -53,7 +46,7 @@ __kernel void m10800_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru
|
|
|
|
|
|
|
|
|
|
sha384_init (&ctx);
|
|
|
|
|
|
|
|
|
|
sha384_update (&ctx, tmp.i, tmp.pw_len);
|
|
|
|
|
sha384_update_swap (&ctx, tmp.i, tmp.pw_len);
|
|
|
|
|
|
|
|
|
|
sha384_final (&ctx);
|
|
|
|
|
|
|
|
|
@ -93,14 +86,7 @@ __kernel void m10800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru
|
|
|
|
|
* base
|
|
|
|
|
*/
|
|
|
|
|
|
|
|
|
|
const u32 pw_len = pws[gid].pw_len;
|
|
|
|
|
|
|
|
|
|
u32 w[64] = { 0 };
|
|
|
|
|
|
|
|
|
|
for (int i = 0, idx = 0; i < pw_len; i += 4, idx += 1)
|
|
|
|
|
{
|
|
|
|
|
w[idx] = swap32_S (pws[gid].i[idx]);
|
|
|
|
|
}
|
|
|
|
|
COPY_PW (pws[gid]);
|
|
|
|
|
|
|
|
|
|
/**
|
|
|
|
|
* loop
|
|
|
|
@ -116,7 +102,7 @@ __kernel void m10800_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru
|
|
|
|
|
|
|
|
|
|
sha384_init (&ctx);
|
|
|
|
|
|
|
|
|
|
sha384_update (&ctx, tmp.i, tmp.pw_len);
|
|
|
|
|
sha384_update_swap (&ctx, tmp.i, tmp.pw_len);
|
|
|
|
|
|
|
|
|
|
sha384_final (&ctx);
|
|
|
|
|
|
|
|
|
|