X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm05800.cl;h=8190a9d978ab7b422477c05f4c8cd117b2c743df;hb=8702d0e3e155e6a61ca1b0dcdc1ec924d92b8436;hp=ad76e9c4b9427a78256484247fe8e5b2cbd64f91;hpb=e4465d7dcb88426532f6b0488cf9bc1cd3b06cc6;p=hashcat.git diff --git a/OpenCL/m05800.cl b/OpenCL/m05800.cl index ad76e9c..8190a9d 100644 --- a/OpenCL/m05800.cl +++ b/OpenCL/m05800.cl @@ -5,20 +5,14 @@ #define _SHA1_ -#include "include/constants.h" -#include "include/kernel_vendor.h" +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" -#define DGST_R0 0 -#define DGST_R1 1 -#define DGST_R2 2 -#define DGST_R3 3 - -#include "include/kernel_functions.c" -#include "types_ocl.c" -#include "common.c" - -#define COMPARE_S "check_single_comp4.c" -#define COMPARE_M "check_multi_comp4.c" +#define COMPARE_S "inc_comp_single.cl" +#define COMPARE_M "inc_comp_multi.cl" typedef struct { @@ -1055,7 +1049,7 @@ __constant entry_t pc[1024] = { 0x33323031, 4 } }; -static void append_word (u32 w0[4], u32 w1[4], const u32 append[4], const u32 offset) +void append_word (u32 w0[4], u32 w1[4], const u32 append[4], const u32 offset) { switch (offset) { @@ -1089,7 +1083,7 @@ static void append_word (u32 w0[4], u32 w1[4], const u32 append[4], const u32 of } } -static void append_salt (u32 w0[4], u32 w1[4], u32 w2[4], const u32 append[5], const u32 offset) +void append_salt (u32 w0[4], u32 w1[4], u32 w2[4], const u32 append[5], const u32 offset) { u32 tmp0; u32 tmp1; @@ -1098,7 +1092,7 @@ static void append_salt (u32 w0[4], u32 w1[4], u32 w2[4], const u32 append[5], c u32 tmp4; u32 tmp5; - #ifdef IS_AMD + #if defined IS_AMD || defined IS_GENERIC const int offset_minus_4 = 4 - (offset & 3); @@ -1180,7 +1174,7 @@ static void append_salt (u32 w0[4], u32 w1[4], u32 w2[4], const u32 append[5], c } } -static void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5]) +void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u32 w3[4], u32 digest[5]) { u32 A = digest[0]; u32 B = digest[1]; @@ -1308,7 +1302,7 @@ static void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], c digest[4] += E; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_init (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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) +__kernel void m05800_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * base @@ -1397,13 +1391,34 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_init (__gl tmps[gid].digest_buf[4] = digest[4]; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_loop (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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) +__kernel void m05800_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_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); + const u32 lid = get_local_id (0); + const u32 lsz = get_local_size (0); + + /** + * cache precomputed conversion table in shared memory + */ + + __local entry_t s_pc[1024]; + + for (u32 i = lid; i < 1024; i += lsz) + { + s_pc[i] = pc[i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + /** + * base + */ u32 word_buf[4]; @@ -1414,8 +1429,6 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_loop (__gl const u32 pw_len = pws[gid].pw_len; - const u32 lid = get_local_id (0); - u32 digest[5]; digest[0] = tmps[gid].digest_buf[0]; @@ -1424,35 +1437,6 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_loop (__gl digest[3] = tmps[gid].digest_buf[3]; digest[4] = tmps[gid].digest_buf[4]; - /** - * cache precomputed conversion table in shared memory - */ - - __local entry_t s_pc[1024]; - - const u32 lid16 = lid * 16; - - s_pc[lid16 + 0] = pc[lid16 + 0]; - s_pc[lid16 + 1] = pc[lid16 + 1]; - s_pc[lid16 + 2] = pc[lid16 + 2]; - s_pc[lid16 + 3] = pc[lid16 + 3]; - s_pc[lid16 + 4] = pc[lid16 + 4]; - s_pc[lid16 + 5] = pc[lid16 + 5]; - s_pc[lid16 + 6] = pc[lid16 + 6]; - s_pc[lid16 + 7] = pc[lid16 + 7]; - s_pc[lid16 + 8] = pc[lid16 + 8]; - s_pc[lid16 + 9] = pc[lid16 + 9]; - s_pc[lid16 + 10] = pc[lid16 + 10]; - s_pc[lid16 + 11] = pc[lid16 + 11]; - s_pc[lid16 + 12] = pc[lid16 + 12]; - s_pc[lid16 + 13] = pc[lid16 + 13]; - s_pc[lid16 + 14] = pc[lid16 + 14]; - s_pc[lid16 + 15] = pc[lid16 + 15]; - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - /** * salt */ @@ -1524,7 +1508,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_loop (__gl tmps[gid].digest_buf[4] = digest[4]; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m05800_comp (__global pw_t *pws, __global gpu_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *esalt_bufs, __global u32 *d_return_buf, __global u32 *d_scryptV_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) +__kernel void m05800_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global androidpin_tmp_t *tmps, __global void *hooks, __global u32 *bitmaps_buf_s1_a, __global u32 *bitmaps_buf_s1_b, __global u32 *bitmaps_buf_s1_c, __global u32 *bitmaps_buf_s1_d, __global u32 *bitmaps_buf_s2_a, __global u32 *bitmaps_buf_s2_b, __global u32 *bitmaps_buf_s2_c, __global u32 *bitmaps_buf_s2_d, __global plain_t *plains_buf, __global digest_t *digests_buf, __global u32 *hashes_shown, __global salt_t *salt_bufs, __global void *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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * modifier