From f5e21aea16697d6730e65189740ae8b7c1361b86 Mon Sep 17 00:00:00 2001 From: DoZ10 Date: Fri, 19 May 2017 08:21:02 -0400 Subject: [PATCH 1/5] Merged with master --- OpenCL/m15500.cl | 109 +++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 109 insertions(+) create mode 100644 OpenCL/m15500.cl diff --git a/OpenCL/m15500.cl b/OpenCL/m15500.cl new file mode 100644 index 000000000..533f54b3b --- /dev/null +++ b/OpenCL/m15500.cl @@ -0,0 +1,109 @@ +/** + * Author......: See docs/credits.txt + * License.....: MIT + */ + +#define NEW_SIMD_CODE + +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" +#include "inc_simd.cl" +#include "inc_rp.h" +#include "inc_rp.cl" + +#define COMPARE_S "inc_comp_single.cl" +#define COMPARE_M "inc_comp_multi.cl" + + +__kernel void m15500_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global argon2_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global argon2_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + if (gid >= gid_max) return; + + u32 w0[4]; + + w0[0] = pws[gid].i[0]; + w0[1] = pws[gid].i[1]; + w0[2] = pws[gid].i[2]; + w0[3] = pws[gid].i[3]; + + u32 w1[4]; + + w1[0] = pws[gid].i[4]; + w1[1] = pws[gid].i[5]; + w1[2] = pws[gid].i[6]; + w1[3] = pws[gid].i[7]; + + 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 pw_len = pws[gid].pw_len; + + /** + * main + */ + + + + /** + * Store tmps + */ + + printf("I was in m15500_init()\n"); +} + +__kernel void m15500_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global argon2_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global argon2_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +{ + /** + * base + */ + + const u32 gid = get_global_id (0); + + if ((gid * VECT_SIZE) >= gid_max) return; + + + /** + * iter1 + */ + + + /** + * Update tmps + */ + + printf("I was in m15500_loop()\n"); + } +} + +__kernel void m15500_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global argon2_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global argon2_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) +{ + const u32 gid = get_global_id (0); + const u32 lid = get_local_id (0); + const u32 lsz = get_local_size (0); + + /** + * Mark hashes + */ + + printf("I was in m15500_comp()\n"); +} From e7dab5d2108e84aa55e9dcd54e881c740874f4a8 Mon Sep 17 00:00:00 2001 From: DoZ10 Date: Fri, 19 May 2017 08:22:18 -0400 Subject: [PATCH 2/5] Removed unused structure --- include/types.h | 17 ----------------- 1 file changed, 17 deletions(-) diff --git a/include/types.h b/include/types.h index a1a036ca5..f892f11a6 100644 --- a/include/types.h +++ b/include/types.h @@ -665,23 +665,6 @@ typedef enum user_options_map * structs */ -typedef struct -{ - u8 digest_length; - u8 key_length; - u8 fanout; - u8 depth; - u32 leaf_length; - u32 node_offset; - u32 xof_length; - u8 node_depth; - u8 inner_length; - u8 reserved[14]; - u8 salt[16]; - u8 personnel[16]; - -} blake2params_t; - typedef struct { u64 h[8]; From 556e64ffe6b0c9df8fc78bc5f933909681c366e6 Mon Sep 17 00:00:00 2001 From: DoZ10 Date: Fri, 19 May 2017 08:25:43 -0400 Subject: [PATCH 3/5] Revert "Removed unused structure" This reverts commit e7dab5d2108e84aa55e9dcd54e881c740874f4a8. --- include/types.h | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/include/types.h b/include/types.h index f892f11a6..a1a036ca5 100644 --- a/include/types.h +++ b/include/types.h @@ -665,6 +665,23 @@ typedef enum user_options_map * structs */ +typedef struct +{ + u8 digest_length; + u8 key_length; + u8 fanout; + u8 depth; + u32 leaf_length; + u32 node_offset; + u32 xof_length; + u8 node_depth; + u8 inner_length; + u8 reserved[14]; + u8 salt[16]; + u8 personnel[16]; + +} blake2params_t; + typedef struct { u64 h[8]; From f75f6120264ea2d3b656842256fa3fa77bca6d7d Mon Sep 17 00:00:00 2001 From: DoZ10 Date: Fri, 19 May 2017 08:26:26 -0400 Subject: [PATCH 4/5] Revert "Merged with master" This reverts commit f5e21aea16697d6730e65189740ae8b7c1361b86. --- OpenCL/m15500.cl | 109 ----------------------------------------------- 1 file changed, 109 deletions(-) delete mode 100644 OpenCL/m15500.cl diff --git a/OpenCL/m15500.cl b/OpenCL/m15500.cl deleted file mode 100644 index 533f54b3b..000000000 --- a/OpenCL/m15500.cl +++ /dev/null @@ -1,109 +0,0 @@ -/** - * Author......: See docs/credits.txt - * License.....: MIT - */ - -#define NEW_SIMD_CODE - -#include "inc_vendor.cl" -#include "inc_hash_constants.h" -#include "inc_hash_functions.cl" -#include "inc_types.cl" -#include "inc_common.cl" -#include "inc_simd.cl" -#include "inc_rp.h" -#include "inc_rp.cl" - -#define COMPARE_S "inc_comp_single.cl" -#define COMPARE_M "inc_comp_multi.cl" - - -__kernel void m15500_init (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global argon2_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global argon2_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) -{ - /** - * base - */ - - const u32 gid = get_global_id (0); - - if (gid >= gid_max) return; - - u32 w0[4]; - - w0[0] = pws[gid].i[0]; - w0[1] = pws[gid].i[1]; - w0[2] = pws[gid].i[2]; - w0[3] = pws[gid].i[3]; - - u32 w1[4]; - - w1[0] = pws[gid].i[4]; - w1[1] = pws[gid].i[5]; - w1[2] = pws[gid].i[6]; - w1[3] = pws[gid].i[7]; - - 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 pw_len = pws[gid].pw_len; - - /** - * main - */ - - - - /** - * Store tmps - */ - - printf("I was in m15500_init()\n"); -} - -__kernel void m15500_loop (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global argon2_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global argon2_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) -{ - /** - * base - */ - - const u32 gid = get_global_id (0); - - if ((gid * VECT_SIZE) >= gid_max) return; - - - /** - * iter1 - */ - - - /** - * Update tmps - */ - - printf("I was in m15500_loop()\n"); - } -} - -__kernel void m15500_comp (__global pw_t *pws, __global const kernel_rule_t *rules_buf, __global const comb_t *combs_buf, __global const bf_t *bfs_buf, __global argon2_tmp_t *tmps, __global void *hooks, __global const u32 *bitmaps_buf_s1_a, __global const u32 *bitmaps_buf_s1_b, __global const u32 *bitmaps_buf_s1_c, __global const u32 *bitmaps_buf_s1_d, __global const u32 *bitmaps_buf_s2_a, __global const u32 *bitmaps_buf_s2_b, __global const u32 *bitmaps_buf_s2_c, __global const u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global const digest_t *digests_buf, __global u32 *hashes_shown, __global const salt_t *salt_bufs, __global argon2_t *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV0_buf, __global u32 *d_scryptV1_buf, __global u32 *d_scryptV2_buf, __global u32 *d_scryptV3_buf, const u32 bitmap_mask, const u32 bitmap_shift1, const u32 bitmap_shift2, const u32 salt_pos, const u32 loop_pos, const u32 loop_cnt, const u32 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) -{ - const u32 gid = get_global_id (0); - const u32 lid = get_local_id (0); - const u32 lsz = get_local_size (0); - - /** - * Mark hashes - */ - - printf("I was in m15500_comp()\n"); -} From 62df9648d6a330459634239c146c178a4514f44b Mon Sep 17 00:00:00 2001 From: DoZ10 Date: Fri, 19 May 2017 08:28:30 -0400 Subject: [PATCH 5/5] Removed unused structure --- include/types.h | 17 ----------------- 1 file changed, 17 deletions(-) diff --git a/include/types.h b/include/types.h index a1a036ca5..f892f11a6 100644 --- a/include/types.h +++ b/include/types.h @@ -665,23 +665,6 @@ typedef enum user_options_map * structs */ -typedef struct -{ - u8 digest_length; - u8 key_length; - u8 fanout; - u8 depth; - u32 leaf_length; - u32 node_offset; - u32 xof_length; - u8 node_depth; - u8 inner_length; - u8 reserved[14]; - u8 salt[16]; - u8 personnel[16]; - -} blake2params_t; - typedef struct { u64 h[8];