From 6534211f36e8aa785da93991cccc91789bcd3be3 Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Thu, 4 Feb 2016 22:20:30 +0100 Subject: [PATCH] Converted to new SIMD: -m 11400 -a 0 --- OpenCL/m11400_a0.cl | 194 +++++++++++++++++--------------------------- OpenCL/m11500_a0.cl | 4 - 2 files changed, 73 insertions(+), 125 deletions(-) diff --git a/OpenCL/m11400_a0.cl b/OpenCL/m11400_a0.cl index ed1faf112..1fbc8987d 100644 --- a/OpenCL/m11400_a0.cl +++ b/OpenCL/m11400_a0.cl @@ -5,6 +5,8 @@ #define _MD5_ +#define NEW_SIMD_CODE + #include "include/constants.h" #include "include/kernel_vendor.h" @@ -18,13 +20,19 @@ #include "OpenCL/common.c" #include "include/rp_kernel.h" #include "OpenCL/rp.c" +#include "OpenCL/simd.c" -#define COMPARE_S "OpenCL/check_single_comp4.c" -#define COMPARE_M "OpenCL/check_multi_comp4.c" +#if VECT_SIZE == 1 +#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i)]) +#elif VECT_SIZE == 2 +#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1]) +#elif VECT_SIZE == 4 +#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3]) +#elif VECT_SIZE == 8 +#define uint_to_hex_lower8(i) (u32x) (l_bin2asc[(i).s0], l_bin2asc[(i).s1], l_bin2asc[(i).s2], l_bin2asc[(i).s3], l_bin2asc[(i).s4], l_bin2asc[(i).s5], l_bin2asc[(i).s6], l_bin2asc[(i).s7]) +#endif -#define uint_to_hex_lower8(i) l_bin2asc[(i)] - -static u32 memcat32 (u32 block0[16], u32 block1[16], const u32 block_len, const u32 append0[4], const u32 append1[4], const u32 append2[4], const u32 append3[4], const u32 append_len) +static u32 memcat32 (u32x block0[16], u32x block1[16], const u32 block_len, const u32x append0[4], const u32x append1[4], const u32x append2[4], const u32x append3[4], const u32 append_len) { const u32 mod = block_len & 3; const u32 div = block_len / 4; @@ -32,35 +40,35 @@ static u32 memcat32 (u32 block0[16], u32 block1[16], const u32 block_len, const #if defined IS_AMD || defined IS_GENERIC const int offset_minus_4 = 4 - mod; - u32 append0_t[4]; + u32x append0_t[4]; append0_t[0] = amd_bytealign (append0[0], 0, offset_minus_4); append0_t[1] = amd_bytealign (append0[1], append0[0], offset_minus_4); append0_t[2] = amd_bytealign (append0[2], append0[1], offset_minus_4); append0_t[3] = amd_bytealign (append0[3], append0[2], offset_minus_4); - u32 append1_t[4]; + u32x append1_t[4]; append1_t[0] = amd_bytealign (append1[0], append0[3], offset_minus_4); append1_t[1] = amd_bytealign (append1[1], append1[0], offset_minus_4); append1_t[2] = amd_bytealign (append1[2], append1[1], offset_minus_4); append1_t[3] = amd_bytealign (append1[3], append1[2], offset_minus_4); - u32 append2_t[4]; + u32x append2_t[4]; append2_t[0] = amd_bytealign (append2[0], append1[3], offset_minus_4); append2_t[1] = amd_bytealign (append2[1], append2[0], offset_minus_4); append2_t[2] = amd_bytealign (append2[2], append2[1], offset_minus_4); append2_t[3] = amd_bytealign (append2[3], append2[2], offset_minus_4); - u32 append3_t[4]; + u32x append3_t[4]; append3_t[0] = amd_bytealign (append3[0], append2[3], offset_minus_4); append3_t[1] = amd_bytealign (append3[1], append3[0], offset_minus_4); append3_t[2] = amd_bytealign (append3[2], append3[1], offset_minus_4); append3_t[3] = amd_bytealign (append3[3], append3[2], offset_minus_4); - u32 append4_t[4]; + u32x append4_t[4]; append4_t[0] = amd_bytealign ( 0, append3[3], offset_minus_4); append4_t[1] = 0; @@ -102,35 +110,35 @@ static u32 memcat32 (u32 block0[16], u32 block1[16], const u32 block_len, const const int selector = (0x76543210 >> (offset_minus_4 * 4)) & 0xffff; - u32 append0_t[4]; + u32x append0_t[4]; append0_t[0] = __byte_perm ( 0, append0[0], selector); append0_t[1] = __byte_perm (append0[0], append0[1], selector); append0_t[2] = __byte_perm (append0[1], append0[2], selector); append0_t[3] = __byte_perm (append0[2], append0[3], selector); - u32 append1_t[4]; + u32x append1_t[4]; append1_t[0] = __byte_perm (append0[3], append1[0], selector); append1_t[1] = __byte_perm (append1[0], append1[1], selector); append1_t[2] = __byte_perm (append1[1], append1[2], selector); append1_t[3] = __byte_perm (append1[2], append1[3], selector); - u32 append2_t[4]; + u32x append2_t[4]; append2_t[0] = __byte_perm (append1[3], append2[0], selector); append2_t[1] = __byte_perm (append2[0], append2[1], selector); append2_t[2] = __byte_perm (append2[1], append2[2], selector); append2_t[3] = __byte_perm (append2[2], append2[3], selector); - u32 append3_t[4]; + u32x append3_t[4]; append3_t[0] = __byte_perm (append2[3], append3[0], selector); append3_t[1] = __byte_perm (append3[0], append3[1], selector); append3_t[2] = __byte_perm (append3[1], append3[2], selector); append3_t[3] = __byte_perm (append3[2], append3[3], selector); - u32 append4_t[4]; + u32x append4_t[4]; append4_t[0] = __byte_perm (append3[3], 0, selector); append4_t[1] = 0; @@ -908,37 +916,14 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, * loop */ - for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < rules_cnt; il_pos += VECT_SIZE) { - u32 w0[4]; + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; - w0[0] = pw_buf0[0]; - w0[1] = pw_buf0[1]; - w0[2] = pw_buf0[2]; - w0[3] = pw_buf0[3]; - - u32 w1[4]; - - w1[0] = pw_buf1[0]; - w1[1] = pw_buf1[1]; - w1[2] = pw_buf1[2]; - w1[3] = pw_buf1[3]; - - u32 w2[4]; - - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - - u32 w3[4]; - - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len); + const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); append_0x80_2x4 (w0, w1, out_len); @@ -950,7 +935,7 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, // append the pass to the salt - u32 block0[16]; + u32x block0[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -969,7 +954,7 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - u32 block1[16]; + u32x block1[16]; block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; @@ -992,28 +977,28 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, block_len = memcat32 (block0, block1, salt_len, w0, w1, w2, w3, out_len); - u32 w0_t[4]; + u32x w0_t[4]; w0_t[0] = block0[ 0]; w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - u32 w1_t[4]; + u32x w1_t[4]; w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - u32 w2_t[4]; + u32x w2_t[4]; w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - u32 w3_t[4]; + u32x w3_t[4]; w3_t[0] = block0[12]; w3_t[1] = block0[13]; @@ -1027,10 +1012,10 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, // md5 - u32 a = MD5M_A; - u32 b = MD5M_B; - u32 c = MD5M_C; - u32 d = MD5M_D; + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00); MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01); @@ -1107,10 +1092,10 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, if (block_len > 55) { - u32 r_a = a; - u32 r_b = b; - u32 r_c = c; - u32 r_d = d; + u32x r_a = a; + u32x r_b = b; + u32x r_c = c; + u32x r_d = d; w0_t[0] = block1[ 0]; w0_t[1] = block1[ 1]; @@ -1319,10 +1304,10 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += MD5M_C; d += MD5M_D; - u32 r_a = a; - u32 r_b = b; - u32 r_c = c; - u32 r_d = d; + u32x r_a = a; + u32x r_b = b; + u32x r_c = c; + u32x r_d = d; // 2nd transform @@ -1533,12 +1518,7 @@ __kernel void m11400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += r_c; d += r_d; - const u32 r0 = a; - const u32 r1 = d; - const u32 r2 = c; - const u32 r3 = b; - - #include COMPARE_M + COMPARE_M_SIMD (a, d, c, b); } } @@ -1725,37 +1705,14 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, * loop */ - for (u32 il_pos = 0; il_pos < rules_cnt; il_pos++) + for (u32 il_pos = 0; il_pos < rules_cnt; il_pos += VECT_SIZE) { - u32 w0[4]; + u32x w0[4] = { 0 }; + u32x w1[4] = { 0 }; + u32x w2[4] = { 0 }; + u32x w3[4] = { 0 }; - w0[0] = pw_buf0[0]; - w0[1] = pw_buf0[1]; - w0[2] = pw_buf0[2]; - w0[3] = pw_buf0[3]; - - u32 w1[4]; - - w1[0] = pw_buf1[0]; - w1[1] = pw_buf1[1]; - w1[2] = pw_buf1[2]; - w1[3] = pw_buf1[3]; - - u32 w2[4]; - - w2[0] = 0; - w2[1] = 0; - w2[2] = 0; - w2[3] = 0; - - u32 w3[4]; - - w3[0] = 0; - w3[1] = 0; - w3[2] = 0; - w3[3] = 0; - - const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len); + const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); append_0x80_2x4 (w0, w1, out_len); @@ -1767,7 +1724,7 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, // append the pass to the salt - u32 block0[16]; + u32x block0[16]; block0[ 0] = salt_buf0[ 0]; block0[ 1] = salt_buf0[ 1]; @@ -1786,7 +1743,7 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, block0[14] = salt_buf0[14]; block0[15] = salt_buf0[15]; - u32 block1[16]; + u32x block1[16]; block1[ 0] = salt_buf1[ 0]; block1[ 1] = salt_buf1[ 1]; @@ -1809,28 +1766,28 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, block_len = memcat32 (block0, block1, salt_len, w0, w1, w2, w3, out_len); - u32 w0_t[4]; + u32x w0_t[4]; w0_t[0] = block0[ 0]; w0_t[1] = block0[ 1]; w0_t[2] = block0[ 2]; w0_t[3] = block0[ 3]; - u32 w1_t[4]; + u32x w1_t[4]; w1_t[0] = block0[ 4]; w1_t[1] = block0[ 5]; w1_t[2] = block0[ 6]; w1_t[3] = block0[ 7]; - u32 w2_t[4]; + u32x w2_t[4]; w2_t[0] = block0[ 8]; w2_t[1] = block0[ 9]; w2_t[2] = block0[10]; w2_t[3] = block0[11]; - u32 w3_t[4]; + u32x w3_t[4]; w3_t[0] = block0[12]; w3_t[1] = block0[13]; @@ -1844,10 +1801,10 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, // md5 - u32 a = MD5M_A; - u32 b = MD5M_B; - u32 c = MD5M_C; - u32 d = MD5M_D; + u32x a = MD5M_A; + u32x b = MD5M_B; + u32x c = MD5M_C; + u32x d = MD5M_D; MD5_STEP (MD5_Fo, a, b, c, d, w0_t[0], MD5C00, MD5S00); MD5_STEP (MD5_Fo, d, a, b, c, w0_t[1], MD5C01, MD5S01); @@ -1924,10 +1881,10 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, if (block_len > 55) { - u32 r_a = a; - u32 r_b = b; - u32 r_c = c; - u32 r_d = d; + u32x r_a = a; + u32x r_b = b; + u32x r_c = c; + u32x r_d = d; w0_t[0] = block1[ 0]; w0_t[1] = block1[ 1]; @@ -2136,10 +2093,10 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += MD5M_C; d += MD5M_D; - u32 r_a = a; - u32 r_b = b; - u32 r_c = c; - u32 r_d = d; + u32x r_a = a; + u32x r_b = b; + u32x r_c = c; + u32x r_d = d; // 2nd transform @@ -2350,12 +2307,7 @@ __kernel void m11400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, c += r_c; d += r_d; - const u32 r0 = a; - const u32 r1 = d; - const u32 r2 = c; - const u32 r3 = b; - - #include COMPARE_S + COMPARE_S_SIMD (a, d, c, b); } } diff --git a/OpenCL/m11500_a0.cl b/OpenCL/m11500_a0.cl index c69dac69d..58aac5468 100644 --- a/OpenCL/m11500_a0.cl +++ b/OpenCL/m11500_a0.cl @@ -182,8 +182,6 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); - append_0x01_2x4 (w0, w1, out_len); - u32x w[16]; w[ 0] = w0[0]; @@ -279,8 +277,6 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, const u32 out_len = apply_rules_vect (pw_buf0, pw_buf1, pw_len, rules_buf, il_pos, w0, w1); - append_0x01_2x4 (w0, w1, out_len); - u32x w[16]; w[ 0] = w0[0];