From 326000035735d24746e5b1baa8fcf2fe7cfb17de Mon Sep 17 00:00:00 2001 From: jsteube Date: Sat, 26 Aug 2017 19:51:37 +0200 Subject: [PATCH] Fix whirlpool pure kernel in -a 0 mode --- OpenCL/m06100_a0.cl | 22 ++++------------------ 1 file changed, 4 insertions(+), 18 deletions(-) diff --git a/OpenCL/m06100_a0.cl b/OpenCL/m06100_a0.cl index 8522ff117..baf8b2c14 100644 --- a/OpenCL/m06100_a0.cl +++ b/OpenCL/m06100_a0.cl @@ -61,14 +61,7 @@ __kernel void m06100_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 @@ -84,7 +77,7 @@ __kernel void m06100_mxx (__global pw_t *pws, __constant const kernel_rule_t *ru whirlpool_init (&ctx, s_Ch, s_Cl); - whirlpool_update (&ctx, tmp.i, tmp.pw_len); + whirlpool_update_swap (&ctx, tmp.i, tmp.pw_len); whirlpool_final (&ctx); @@ -155,14 +148,7 @@ __kernel void m06100_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 @@ -178,7 +164,7 @@ __kernel void m06100_sxx (__global pw_t *pws, __constant const kernel_rule_t *ru whirlpool_init (&ctx, s_Ch, s_Cl); - whirlpool_update (&ctx, tmp.i, tmp.pw_len); + whirlpool_update_swap (&ctx, tmp.i, tmp.pw_len); whirlpool_final (&ctx);