From c0a38846e3223ebc404b7dd7eb0f7e64dff86fe9 Mon Sep 17 00:00:00 2001 From: jsteube Date: Tue, 3 May 2016 10:14:53 +0200 Subject: [PATCH] Fix -m 5100 multihash and benchmark --- OpenCL/m05100_a0.cl | 38 ++++++++++---------------------------- OpenCL/m05100_a1.cl | 26 ++++++++++---------------- OpenCL/m05100_a3.cl | 26 ++++++++++---------------- 3 files changed, 30 insertions(+), 60 deletions(-) diff --git a/OpenCL/m05100_a0.cl b/OpenCL/m05100_a0.cl index 9f99183..5f6593b 100644 --- a/OpenCL/m05100_a0.cl +++ b/OpenCL/m05100_a0.cl @@ -52,18 +52,6 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, const u32 pw_len = pws[gid].pw_len; - /** - * digest - */ - - const u32 search[4] = - { - digests_buf[digests_offset].digest_buf[DGST_R0], - digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] - }; - /** * loop */ @@ -164,14 +152,11 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += MD5M_C; d += MD5M_D; - u32x e = 0; - u32x f = 0; - - COMPARE_M_SIMD (a, b, e, f); + u32x z = 0; - COMPARE_M_SIMD (b, c, e, f); - - COMPARE_M_SIMD (c, d, e, f); + COMPARE_M_SIMD (a, b, z, z); + COMPARE_M_SIMD (b, c, z, z); + COMPARE_M_SIMD (c, d, z, z); } } @@ -221,8 +206,8 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** @@ -325,14 +310,11 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += MD5M_C; d += MD5M_D; - u32x e = 0; - u32x f = 0; - - COMPARE_S_SIMD (a, b, e, f); - - COMPARE_S_SIMD (b, c, e, f); + u32x z = 0; - COMPARE_S_SIMD (c, d, e, f); + COMPARE_S_SIMD (a, b, z, z); + COMPARE_S_SIMD (b, c, z, z); + COMPARE_S_SIMD (c, d, z, z); } } diff --git a/OpenCL/m05100_a1.cl b/OpenCL/m05100_a1.cl index d9eb91d..795f883 100644 --- a/OpenCL/m05100_a1.cl +++ b/OpenCL/m05100_a1.cl @@ -205,14 +205,11 @@ __kernel void m05100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += MD5M_C; d += MD5M_D; - u32x e = 0; - u32x f = 0; + u32x z = 0; - COMPARE_M_SIMD (a, b, e, f); - - COMPARE_M_SIMD (b, c, e, f); - - COMPARE_M_SIMD (c, d, e, f); + COMPARE_M_SIMD (a, b, z, z); + COMPARE_M_SIMD (b, c, z, z); + COMPARE_M_SIMD (c, d, z, z); } } @@ -262,8 +259,8 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** @@ -421,14 +418,11 @@ __kernel void m05100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += MD5M_C; d += MD5M_D; - u32x e = 0; - u32x f = 0; - - COMPARE_S_SIMD (a, b, e, f); - - COMPARE_S_SIMD (b, c, e, f); + u32x z = 0; - COMPARE_S_SIMD (c, d, e, f); + COMPARE_S_SIMD (a, b, z, z); + COMPARE_S_SIMD (b, c, z, z); + COMPARE_S_SIMD (c, d, z, z); } } diff --git a/OpenCL/m05100_a3.cl b/OpenCL/m05100_a3.cl index 90365c9..5a41ae2 100644 --- a/OpenCL/m05100_a3.cl +++ b/OpenCL/m05100_a3.cl @@ -145,14 +145,11 @@ void m05100m (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl c += MD5M_C; d += MD5M_D; - u32x e = 0; - u32x f = 0; + u32x z = 0; - COMPARE_M_SIMD (a, b, e, f); - - COMPARE_M_SIMD (b, c, e, f); - - COMPARE_M_SIMD (c, d, e, f); + COMPARE_M_SIMD (a, b, z, z); + COMPARE_M_SIMD (b, c, z, z); + COMPARE_M_SIMD (c, d, z, z); } } @@ -173,8 +170,8 @@ void m05100s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl { digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** @@ -293,14 +290,11 @@ void m05100s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl c += MD5M_C; d += MD5M_D; - u32x e = 0; - u32x f = 0; - - COMPARE_S_SIMD (a, b, e, f); - - COMPARE_S_SIMD (b, c, e, f); + u32x z = 0; - COMPARE_S_SIMD (c, d, e, f); + COMPARE_S_SIMD (a, b, z, z); + COMPARE_S_SIMD (b, c, z, z); + COMPARE_S_SIMD (c, d, z, z); } } -- 2.25.1