X-Git-Url: https://www.flypig.org.uk/git/?a=blobdiff_plain;f=OpenCL%2Fm06232.cl;h=cea0016253afd42c33dc7a35e9a084b5ac423634;hb=161a6eb4bc643d8e636e96eda613f5137d30da59;hp=96c0579a3623aaa6126b79f1535e443cf6f8bfb9;hpb=083daf4613259e9c3d52530386a92db76963af92;p=hashcat.git diff --git a/OpenCL/m06232.cl b/OpenCL/m06232.cl index 96c0579..cea0016 100644 --- a/OpenCL/m06232.cl +++ b/OpenCL/m06232.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/gpu_aes256_amd.c" -#include "OpenCL/gpu_twofish256_amd.c" -#include "OpenCL/gpu_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 (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] = { @@ -1170,7 +1166,9 @@ static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch u32 i; - #pragma unroll 8 + #ifdef _unroll + #pragma unroll + #endif for (i = 0; i < 8; i++) { const u8 Lp0 = Kh[(i + 8) & 7] >> 24; @@ -1218,7 +1216,9 @@ static void whirlpool_transform (const u32 w[16], u32 dgst[16], __local u32 s_Ch Kh[7] = Lh[7]; Kl[7] = Ll[7]; - #pragma unroll 8 + #ifdef _unroll + #pragma unroll + #endif for (i = 0; i < 8; i++) { const u8 Lp0 = stateh[(i + 8) & 7] >> 24; @@ -1285,7 +1285,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_run2 (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_run2 (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]; @@ -1365,7 +1365,7 @@ static void hmac_run2 (const u32 w1[16], const u32 w2[16], const u32 ipad[16], c whirlpool_transform (w, 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; @@ -1440,7 +1440,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; @@ -1465,13 +1465,51 @@ static u32 u8add (const u32 a, const u32 b) return r; } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_init (__global pw_t *pws, __global gpu_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 m06232_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]; @@ -1522,33 +1560,6 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_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 */ @@ -1695,34 +1706,50 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_init (__gl } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_loop (__global pw_t *pws, __global gpu_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 m06232_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 + */ - for (u32 i = 0; i < 8; i++) + __local u32 s_Ch[8][256]; + __local u32 s_Cl[8][256]; + + 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]; @@ -1897,7 +1924,7 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_loop (__gl } } -__kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_comp (__global pw_t *pws, __global gpu_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 m06232_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 @@ -1937,6 +1964,8 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_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]; { @@ -1947,11 +1976,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_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); } } @@ -1963,11 +1990,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_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); } } @@ -1979,11 +2004,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_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); } } @@ -2018,11 +2041,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_comp (__gl aes256_decrypt_xts (ukey2, ukey4, tmp, tmp); twofish256_decrypt_xts (ukey1, ukey3, 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); } } @@ -2035,11 +2056,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_comp (__gl serpent256_decrypt_xts (ukey2, ukey4, tmp, tmp); aes256_decrypt_xts (ukey1, ukey3, 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); } } @@ -2052,11 +2071,9 @@ __kernel void __attribute__((reqd_work_group_size (64, 1, 1))) m06232_comp (__gl twofish256_decrypt_xts (ukey2, ukey4, tmp, tmp); serpent256_decrypt_xts (ukey1, ukey3, 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); } } }