X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm06231.cl;h=4bcf4458fc6bd70a8533b34d1ca6ec89460c1b0d;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=432a780af0cf421c3cb225b47fc3fdbebf12f54c;hpb=331188167cc7b4f6f2a336d63daca23fbdfc199c;p=hashcat.git diff --git a/OpenCL/m06231.cl b/OpenCL/m06231.cl index 432a780..4bcf445 100644 --- a/OpenCL/m06231.cl +++ b/OpenCL/m06231.cl @@ -1,25 +1,21 @@ /** - * Author......: Jens Steube + * Authors.....: Jens Steube + * Gabriele Gristina + * * License.....: MIT */ #define _WHIRLPOOL_ -#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 "OpenCL/types_ocl.c" -#include "OpenCL/common.c" - -#include "OpenCL/kernel_aes256_amd.c" -#include "OpenCL/kernel_twofish256_amd.c" -#include "OpenCL/kernel_serpent256_amd.c" +#include "inc_cipher_aes256.cl" +#include "inc_cipher_twofish256.cl" +#include "inc_cipher_serpent256.cl" #define R 10 @@ -1089,7 +1085,7 @@ __constant u32 Cl[8][256] = #define BOX(S,n,i) (S)[(n)][(i)] -static void whirlpool_transform_last (u32 dgst[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256]) +void whirlpool_transform_last (u32 dgst[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256]) { const u32 rch[R + 1] = { @@ -1173,6 +1169,9 @@ static void whirlpool_transform_last (u32 dgst[16], __local u32 s_Ch[8][256], __ u32 i; + #ifdef _unroll + #pragma unroll + #endif for (i = 0; i < 8; i++) { const u32 Lp0 = Kh[(i + 8) & 7] >> 24; @@ -1220,6 +1219,9 @@ static void whirlpool_transform_last (u32 dgst[16], __local u32 s_Ch[8][256], __ Kh[7] = Lh[7]; Kl[7] = Ll[7]; + #ifdef _unroll + #pragma unroll + #endif for (i = 0; i < 8; i++) { const u32 Lp0 = stateh[(i + 8) & 7] >> 24; @@ -1286,7 +1288,7 @@ static void whirlpool_transform_last (u32 dgst[16], __local u32 s_Ch[8][256], __ dgst[15] ^= statel[7] ^ LAST_W15; } -static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256]) +void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256]) { const u32 rch[R + 1] = { @@ -1367,6 +1369,9 @@ static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch u32 i; + #ifdef _unroll + #pragma unroll + #endif for (i = 0; i < 8; i++) { const u32 Lp0 = Kh[(i + 8) & 7] >> 24; @@ -1414,6 +1419,9 @@ static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch Kh[7] = Lh[7]; Kl[7] = Ll[7]; + #ifdef _unroll + #pragma unroll + #endif for (i = 0; i < 8; i++) { const u32 Lp0 = stateh[(i + 8) & 7] >> 24; @@ -1480,7 +1488,7 @@ static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch dgst[15] ^= statel[7] ^ w[15]; } -static void hmac_run2a (const u32 w1[16], const u32 w2[16], const u32 ipad[16], const u32 opad[16], u32 dgst[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256]) +void hmac_run2a (const u32 w1[16], const u32 w2[16], const u32 ipad[16], const u32 opad[16], u32 dgst[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256]) { dgst[ 0] = ipad[ 0]; dgst[ 1] = ipad[ 1]; @@ -1543,7 +1551,7 @@ static void hmac_run2a (const u32 w1[16], const u32 w2[16], const u32 ipad[16], whirlpool_transform_last (dgst, s_Ch, s_Cl); } -static void hmac_run2b (const u32 w1[16], const u32 ipad[16], const u32 opad[16], u32 dgst[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256]) +void hmac_run2b (const u32 w1[16], const u32 ipad[16], const u32 opad[16], u32 dgst[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256]) { dgst[ 0] = ipad[ 0]; dgst[ 1] = ipad[ 1]; @@ -1607,7 +1615,7 @@ static void hmac_run2b (const u32 w1[16], const u32 ipad[16], const u32 opad[16] whirlpool_transform_last (dgst, s_Ch, s_Cl); } -static void hmac_init (u32 w[16], u32 ipad[16], u32 opad[16], __local u32 s_Ch[8][256], __local u32 s_Cl[8][256]) +void hmac_init (u32 w[16], u32 ipad[16], u32 opad[16], __local u32 (*s_Ch)[256], __local u32 (*s_Cl)[256]) { w[ 0] ^= 0x36363636; w[ 1] ^= 0x36363636; @@ -1682,7 +1690,7 @@ static void hmac_init (u32 w[16], u32 ipad[16], u32 opad[16], __local u32 s_Ch[8 whirlpool_transform (w, opad, s_Ch, s_Cl); } -static u32 u8add (const u32 a, const u32 b) +u32 u8add (const u32 a, const u32 b) { const u32 a1 = (a >> 0) & 0xff; const u32 a2 = (a >> 8) & 0xff; @@ -1707,13 +1715,51 @@ static u32 u8add (const u32 a, const u32 b) return r; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_t *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 m06231_init (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** - * base + * modifier */ const u32 gid = get_global_id (0); + const u32 lid = get_local_id (0); + const u32 lsz = get_local_size (0); + + /** + * shared + */ + + __local u32 s_Ch[8][256]; + __local u32 s_Cl[8][256]; + + for (u32 i = lid; i < 256; i += lsz) + { + s_Ch[0][i] = Ch[0][i]; + s_Ch[1][i] = Ch[1][i]; + s_Ch[2][i] = Ch[2][i]; + s_Ch[3][i] = Ch[3][i]; + s_Ch[4][i] = Ch[4][i]; + s_Ch[5][i] = Ch[5][i]; + s_Ch[6][i] = Ch[6][i]; + s_Ch[7][i] = Ch[7][i]; + + s_Cl[0][i] = Cl[0][i]; + s_Cl[1][i] = Cl[1][i]; + s_Cl[2][i] = Cl[2][i]; + s_Cl[3][i] = Cl[3][i]; + s_Cl[4][i] = Cl[4][i]; + s_Cl[5][i] = Cl[5][i]; + s_Cl[6][i] = Cl[6][i]; + s_Cl[7][i] = Cl[7][i]; + } + + barrier (CLK_LOCAL_MEM_FENCE); + + if (gid >= gid_max) return; + + /** + * base + */ u32 w0[4]; @@ -1764,33 +1810,6 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_init (__gl w3[2] = u8add (w3[2], esalt_bufs[salt_pos].keyfile_buf[14]); w3[3] = u8add (w3[3], esalt_bufs[salt_pos].keyfile_buf[15]); - /** - * shared mem - */ - - __local u32 s_Ch[8][256]; - __local u32 s_Cl[8][256]; - - const u32 lid = get_local_id (0); - - const u32 lid4 = lid * 4; - - for (u32 i = 0; i < 8; i++) - { - s_Ch[i][lid4 + 0] = Ch[i][lid4 + 0]; - s_Ch[i][lid4 + 1] = Ch[i][lid4 + 1]; - s_Ch[i][lid4 + 2] = Ch[i][lid4 + 2]; - s_Ch[i][lid4 + 3] = Ch[i][lid4 + 3]; - s_Cl[i][lid4 + 0] = Cl[i][lid4 + 0]; - s_Cl[i][lid4 + 1] = Cl[i][lid4 + 1]; - s_Cl[i][lid4 + 2] = Cl[i][lid4 + 2]; - s_Cl[i][lid4 + 3] = Cl[i][lid4 + 3]; - } - - barrier (CLK_LOCAL_MEM_FENCE); - - if (gid >= gid_max) return; - /** * salt */ @@ -1937,34 +1956,50 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_init (__gl } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_t *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 m06231_loop (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - - __local u32 s_Ch[8][256]; - __local u32 s_Cl[8][256]; + /** + * modifier + */ const u32 gid = get_global_id (0); const u32 lid = get_local_id (0); + const u32 lsz = get_local_size (0); - const u32 lid4 = lid * 4; + /** + * shared + */ + + __local u32 s_Ch[8][256]; + __local u32 s_Cl[8][256]; - for (u32 i = 0; i < 8; i++) + for (u32 i = lid; i < 256; i += lsz) { - s_Ch[i][lid4 + 0] = Ch[i][lid4 + 0]; - s_Ch[i][lid4 + 1] = Ch[i][lid4 + 1]; - s_Ch[i][lid4 + 2] = Ch[i][lid4 + 2]; - s_Ch[i][lid4 + 3] = Ch[i][lid4 + 3]; - s_Cl[i][lid4 + 0] = Cl[i][lid4 + 0]; - s_Cl[i][lid4 + 1] = Cl[i][lid4 + 1]; - s_Cl[i][lid4 + 2] = Cl[i][lid4 + 2]; - s_Cl[i][lid4 + 3] = Cl[i][lid4 + 3]; + s_Ch[0][i] = Ch[0][i]; + s_Ch[1][i] = Ch[1][i]; + s_Ch[2][i] = Ch[2][i]; + s_Ch[3][i] = Ch[3][i]; + s_Ch[4][i] = Ch[4][i]; + s_Ch[5][i] = Ch[5][i]; + s_Ch[6][i] = Ch[6][i]; + s_Ch[7][i] = Ch[7][i]; + + s_Cl[0][i] = Cl[0][i]; + s_Cl[1][i] = Cl[1][i]; + s_Cl[2][i] = Cl[2][i]; + s_Cl[3][i] = Cl[3][i]; + s_Cl[4][i] = Cl[4][i]; + s_Cl[5][i] = Cl[5][i]; + s_Cl[6][i] = Cl[6][i]; + s_Cl[7][i] = Cl[7][i]; } barrier (CLK_LOCAL_MEM_FENCE); if (gid >= gid_max) return; + const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; + u32 ipad[16]; ipad[ 0] = tmps[gid].ipad[ 0]; @@ -2120,7 +2155,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_loop (__gl } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_t *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 m06231_comp (__global pw_t *pws, __global kernel_rule_t *rules_buf, __global comb_t *combs_buf, __global bf_t *bfs_buf, __global tc_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 tc_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 il_cnt, const u32 digests_cnt, const u32 digests_offset, const u32 combs_mode, const u32 gid_max) { /** * base @@ -2160,6 +2195,8 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_comp (__gl data[2] = esalt_bufs[0].data_buf[2]; data[3] = esalt_bufs[0].data_buf[3]; + const u32 signature = esalt_bufs[0].signature; + u32 tmp[4]; { @@ -2170,11 +2207,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_comp (__gl aes256_decrypt_xts (ukey1, ukey2, tmp, tmp); - if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5))) + if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); - - d_return_buf[lid] = 1; + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); } } @@ -2186,11 +2221,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_comp (__gl serpent256_decrypt_xts (ukey1, ukey2, tmp, tmp); - if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5))) + if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); - - d_return_buf[lid] = 1; + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); } } @@ -2202,11 +2235,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06231_comp (__gl twofish256_decrypt_xts (ukey1, ukey2, tmp, tmp); - if (((tmp[0] == 0x45555254) && (tmp[3] == 0)) || ((tmp[0] == 0x45555254) && ((tmp[1] >> 16) <= 5))) + if (((tmp[0] == signature) && (tmp[3] == 0)) || ((tmp[0] == signature) && ((tmp[1] >> 16) <= 5))) { - mark_hash (plains_buf, hashes_shown, 0, gid, 0); - - d_return_buf[lid] = 1; + mark_hash (plains_buf, d_return_buf, salt_pos, 0, 0, gid, 0); } } }