X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm13200.cl;h=2b05fc47786154bace53e3150127ae056a967878;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=c0b4af5e6d1ff95d107925c8e4d88f1ecf27adea;hpb=a7886f02f44ebc3d01f2c9ee3fef0c35e2736b96;p=hashcat.git diff --git a/OpenCL/m13200.cl b/OpenCL/m13200.cl index c0b4af5..2b05fc4 100644 --- a/OpenCL/m13200.cl +++ b/OpenCL/m13200.cl @@ -4,20 +4,14 @@ * License.....: MIT */ - -#define _SHA1_ -#include "include/constants.h" -#include "include/kernel_vendor.h" +#define _AXCRYPT_ -#define DGST_R0 0 -#define DGST_R1 1 -#define DGST_R2 2 -#define DGST_R3 3 - -#include "include/kernel_functions.c" -#include "OpenCL/types_ocl.c" -#include "OpenCL/common.c" +#include "inc_vendor.cl" +#include "inc_hash_constants.h" +#include "inc_hash_functions.cl" +#include "inc_types.cl" +#include "inc_common.cl" __constant u32 te0[256] = { @@ -713,7 +707,9 @@ void AES128_ExpandKey (u32 *userkey, u32 *rek, __local u32 *s_te0, __local u32 * rek[2] = userkey[2]; rek[3] = userkey[3]; - #pragma unroll 10 + #ifdef _unroll + #pragma unroll + #endif for (u32 i = 0, j = 0; i < 10; i += 1, j += 4) { u32 temp = rek[j + 3]; @@ -975,7 +971,7 @@ void sha1_transform (const u32 w0[4], const u32 w1[4], const u32 w2[4], const u3 digest[4] += E; } -__kernel void m13200_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global axcrypt_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 m13200_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global axcrypt_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * base @@ -986,28 +982,22 @@ __kernel void m13200_init (__global pw_t *pws, __global kernel_rule_t *rules_buf if (gid >= gid_max) return; u32 w0[4]; + u32 w1[4]; + u32 w2[4]; + u32 w3[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] = pws[gid].i[ 8]; w2[1] = pws[gid].i[ 9]; w2[2] = pws[gid].i[10]; w2[3] = pws[gid].i[11]; - - u32 w3[4]; - w3[0] = pws[gid].i[12]; w3[1] = pws[gid].i[13]; w3[2] = pws[gid].i[14]; @@ -1016,7 +1006,7 @@ __kernel void m13200_init (__global pw_t *pws, __global kernel_rule_t *rules_buf const u32 pw_len = pws[gid].pw_len; append_0x80_4x4 (w0, w1, w2, w3, pw_len); - + w0[0] = swap32 (w0[0]); w0[1] = swap32 (w0[1]); w0[2] = swap32 (w0[2]); @@ -1031,11 +1021,9 @@ __kernel void m13200_init (__global pw_t *pws, __global kernel_rule_t *rules_buf w2[3] = swap32 (w2[3]); w3[0] = swap32 (w3[0]); w3[1] = swap32 (w3[1]); - w3[2] = swap32 (w3[2]); - w3[3] = swap32 (w3[3]); - + w3[2] = 0; w3[3] = pw_len * 8; - + /** * KEK */ @@ -1072,10 +1060,10 @@ __kernel void m13200_init (__global pw_t *pws, __global kernel_rule_t *rules_buf tmps[gid].cipher[1] = salt_bufs[salt_pos].salt_buf[5]; tmps[gid].cipher[2] = 0; tmps[gid].cipher[3] = 0; - + } -__kernel void m13200_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global axcrypt_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 m13200_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global axcrypt_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 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); @@ -1137,17 +1125,19 @@ __kernel void m13200_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf ukeyx[1] = tmps[gid].KEK[1]; ukeyx[2] = tmps[gid].KEK[2]; ukeyx[3] = tmps[gid].KEK[3]; - + AES128_ExpandKey (ukeyx, rek, s_te0, s_te1, s_te2, s_te3, s_te4); - #pragma unroll KEYLEN + #ifdef _unroll + #pragma unroll + #endif for (u32 i = 0; i < KEYLEN; i++) rdk[i] = rek[i]; AES128_InvertKey (rdk, s_td0, s_td1, s_td2, s_td3, s_td4, s_te0, s_te1, s_te2, s_te3, s_te4); u32 lsb[4]; u32 cipher[4]; - + lsb[0] = tmps[gid].lsb[0]; lsb[1] = tmps[gid].lsb[1]; lsb[2] = tmps[gid].lsb[2]; @@ -1157,8 +1147,8 @@ __kernel void m13200_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf cipher[1] = tmps[gid].cipher[1]; cipher[2] = tmps[gid].cipher[2]; cipher[3] = tmps[gid].cipher[3]; - - + + /** * AxCrypt main cipher routine */ @@ -1201,14 +1191,14 @@ __kernel void m13200_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf tmps[gid].lsb[1] = lsb[1]; tmps[gid].lsb[2] = lsb[2]; tmps[gid].lsb[3] = lsb[3]; - + tmps[gid].cipher[0] = cipher[0]; tmps[gid].cipher[1] = cipher[1]; tmps[gid].cipher[2] = cipher[2]; tmps[gid].cipher[3] = cipher[3]; } -__kernel void m13200_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global axcrypt_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 m13200_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global axcrypt_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 rules_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * base @@ -1223,8 +1213,6 @@ __kernel void m13200_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf if(tmps[gid].cipher[0]==0xA6A6A6A6 && tmps[gid].cipher[1]==0xA6A6A6A6) { - mark_hash (plains_buf, hashes_shown, digests_offset, gid, il_pos); - - d_return_buf[lid] = 1; + mark_hash (plains_buf, d_return_buf, salt_pos, 0, digests_offset + 0, gid, il_pos); } }