From 441434840ceb686eb4e589bd17757bf2b5b0ddb3 Mon Sep 17 00:00:00 2001 From: jsteube Date: Wed, 19 Jul 2017 17:45:58 +0200 Subject: [PATCH] Fix broken -m 7900 after migration to pure kernel --- OpenCL/m07900.cl | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/OpenCL/m07900.cl b/OpenCL/m07900.cl index d49b9d71b..d50b1d4c2 100644 --- a/OpenCL/m07900.cl +++ b/OpenCL/m07900.cl @@ -194,10 +194,10 @@ __kernel void m07900_loop (__global pw_t *pws, __global const kernel_rule_t *rul digest[1] = SHA512M_B; digest[2] = SHA512M_C; digest[3] = SHA512M_D; - digest[4] = SHA512M_A; - digest[5] = SHA512M_B; - digest[6] = SHA512M_C; - digest[7] = SHA512M_D; + digest[4] = SHA512M_E; + digest[5] = SHA512M_F; + digest[6] = SHA512M_G; + digest[7] = SHA512M_H; sha512_transform (sha512_ctx.w0, sha512_ctx.w1, sha512_ctx.w2, sha512_ctx.w3, sha512_ctx.w4, sha512_ctx.w5, sha512_ctx.w6, sha512_ctx.w7, digest); } @@ -229,10 +229,10 @@ __kernel void m07900_comp (__global pw_t *pws, __global const kernel_rule_t *rul * digest */ - const u32 r0 = l32_from_64 (tmps[gid].digest_buf[0]); - const u32 r1 = h32_from_64 (tmps[gid].digest_buf[0]); - const u32 r2 = l32_from_64 (tmps[gid].digest_buf[1]); - const u32 r3 = h32_from_64 (tmps[gid].digest_buf[1]); + const u32 r0 = l32_from_64_S (tmps[gid].digest_buf[0]); + const u32 r1 = h32_from_64_S (tmps[gid].digest_buf[0]); + const u32 r2 = l32_from_64_S (tmps[gid].digest_buf[1]); + const u32 r3 = h32_from_64_S (tmps[gid].digest_buf[1]); #define il_pos 0