From cad3b3e10bf7f66a2e0311b5967b7d4722e93569 Mon Sep 17 00:00:00 2001 From: jsteube Date: Wed, 30 Nov 2016 11:41:31 +0100 Subject: [PATCH] Get rid of truecrypt_mdlen --- OpenCL/inc_types.cl | 1 - OpenCL/m06211.cl | 8 ++------ OpenCL/m06212.cl | 8 ++------ OpenCL/m06213.cl | 8 ++------ OpenCL/m06221.cl | 8 ++------ OpenCL/m06222.cl | 8 ++------ OpenCL/m06223.cl | 8 ++------ OpenCL/m06231.cl | 8 ++------ OpenCL/m06232.cl | 8 ++------ OpenCL/m06233.cl | 8 ++------ OpenCL/m13751.cl | 8 ++------ OpenCL/m13752.cl | 8 ++------ OpenCL/m13753.cl | 8 ++------ include/types.h | 1 - src/interface.c | 50 ++++----------------------------------------- 15 files changed, 28 insertions(+), 120 deletions(-) diff --git a/OpenCL/inc_types.cl b/OpenCL/inc_types.cl index 752acd63c..408481b31 100644 --- a/OpenCL/inc_types.cl +++ b/OpenCL/inc_types.cl @@ -677,7 +677,6 @@ typedef struct u32 salt_sign[2]; u32 keccak_mdlen; - u32 truecrypt_mdlen; u32 digests_cnt; u32 digests_done; diff --git a/OpenCL/m06211.cl b/OpenCL/m06211.cl index d9763b851..d532bfe23 100644 --- a/OpenCL/m06211.cl +++ b/OpenCL/m06211.cl @@ -466,8 +466,6 @@ __kernel void m06211_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf2[14] = (64 + 64 + 4) * 8; salt_buf2[15] = 0; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 w[16]; w[ 0] = w0[0]; @@ -504,7 +502,7 @@ __kernel void m06211_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[3] = opad[3]; tmps[gid].opad[4] = opad[4]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 5, j += 1) + for (u32 i = 0, j = 1; i < 16; i += 5, j += 1) { salt_buf2[0] = swap32 (j); @@ -528,8 +526,6 @@ __kernel void m06211_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m06211_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 tc_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -549,7 +545,7 @@ __kernel void m06211_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[3] = tmps[gid].opad[3]; opad[4] = tmps[gid].opad[4]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 5) + for (u32 i = 0; i < 16; i += 5) { u32 dgst[5]; u32 out[5]; diff --git a/OpenCL/m06212.cl b/OpenCL/m06212.cl index 32cb4df2a..8b5dff286 100644 --- a/OpenCL/m06212.cl +++ b/OpenCL/m06212.cl @@ -466,8 +466,6 @@ __kernel void m06212_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf2[14] = (64 + 64 + 4) * 8; salt_buf2[15] = 0; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 w[16]; w[ 0] = w0[0]; @@ -504,7 +502,7 @@ __kernel void m06212_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[3] = opad[3]; tmps[gid].opad[4] = opad[4]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 5, j += 1) + for (u32 i = 0, j = 1; i < 32; i += 5, j += 1) { salt_buf2[0] = swap32 (j); @@ -528,8 +526,6 @@ __kernel void m06212_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m06212_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 tc_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -549,7 +545,7 @@ __kernel void m06212_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[3] = tmps[gid].opad[3]; opad[4] = tmps[gid].opad[4]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 5) + for (u32 i = 0; i < 32; i += 5) { u32 dgst[5]; u32 out[5]; diff --git a/OpenCL/m06213.cl b/OpenCL/m06213.cl index 58e55d9cf..8ca2bcc66 100644 --- a/OpenCL/m06213.cl +++ b/OpenCL/m06213.cl @@ -466,8 +466,6 @@ __kernel void m06213_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf2[14] = (64 + 64 + 4) * 8; salt_buf2[15] = 0; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 w[16]; w[ 0] = w0[0]; @@ -504,7 +502,7 @@ __kernel void m06213_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[3] = opad[3]; tmps[gid].opad[4] = opad[4]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 5, j += 1) + for (u32 i = 0, j = 1; i < 48; i += 5, j += 1) { salt_buf2[0] = swap32 (j); @@ -528,8 +526,6 @@ __kernel void m06213_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m06213_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 tc_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -549,7 +545,7 @@ __kernel void m06213_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[3] = tmps[gid].opad[3]; opad[4] = tmps[gid].opad[4]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 5) + for (u32 i = 0; i < 48; i += 5) { u32 dgst[5]; u32 out[5]; diff --git a/OpenCL/m06221.cl b/OpenCL/m06221.cl index 53082f408..48d5f908c 100644 --- a/OpenCL/m06221.cl +++ b/OpenCL/m06221.cl @@ -339,8 +339,6 @@ __kernel void m06221_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf[14] = 0; salt_buf[15] = (128 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u64 w[16]; w[ 0] = ((u64) swap32 (w0[0])) << 32 | (u64) swap32 (w0[1]); @@ -383,7 +381,7 @@ __kernel void m06221_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[6] = opad[6]; tmps[gid].opad[7] = opad[7]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 8); i += 8, j += 1) + for (u32 i = 0, j = 1; i < 8; i += 8, j += 1) { salt_buf[8] = (u64) j << 32 | (u64) 0x80000000; @@ -413,8 +411,6 @@ __kernel void m06221_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m06221_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 tc64_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -441,7 +437,7 @@ __kernel void m06221_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[6] = tmps[gid].opad[6]; opad[7] = tmps[gid].opad[7]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 8); i += 8) + for (u32 i = 0; i < 8; i += 8) { u64 dgst[8]; diff --git a/OpenCL/m06222.cl b/OpenCL/m06222.cl index cf11b202c..d4095d561 100644 --- a/OpenCL/m06222.cl +++ b/OpenCL/m06222.cl @@ -339,8 +339,6 @@ __kernel void m06222_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf[14] = 0; salt_buf[15] = (128 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u64 w[16]; w[ 0] = ((u64) swap32 (w0[0])) << 32 | (u64) swap32 (w0[1]); @@ -383,7 +381,7 @@ __kernel void m06222_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[6] = opad[6]; tmps[gid].opad[7] = opad[7]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 8); i += 8, j += 1) + for (u32 i = 0, j = 1; i < 16; i += 8, j += 1) { salt_buf[8] = (u64) j << 32 | (u64) 0x80000000; @@ -413,8 +411,6 @@ __kernel void m06222_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m06222_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 tc64_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -441,7 +437,7 @@ __kernel void m06222_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[6] = tmps[gid].opad[6]; opad[7] = tmps[gid].opad[7]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 8); i += 8) + for (u32 i = 0; i < 16; i += 8) { u64 dgst[8]; diff --git a/OpenCL/m06223.cl b/OpenCL/m06223.cl index 43ac1428f..9618c866a 100644 --- a/OpenCL/m06223.cl +++ b/OpenCL/m06223.cl @@ -339,8 +339,6 @@ __kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf[14] = 0; salt_buf[15] = (128 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u64 w[16]; w[ 0] = ((u64) swap32 (w0[0])) << 32 | (u64) swap32 (w0[1]); @@ -383,7 +381,7 @@ __kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[6] = opad[6]; tmps[gid].opad[7] = opad[7]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 8); i += 8, j += 1) + for (u32 i = 0, j = 1; i < 24; i += 8, j += 1) { salt_buf[8] = (u64) j << 32 | (u64) 0x80000000; @@ -413,8 +411,6 @@ __kernel void m06223_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m06223_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 tc64_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -441,7 +437,7 @@ __kernel void m06223_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[6] = tmps[gid].opad[6]; opad[7] = tmps[gid].opad[7]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 8); i += 8) + for (u32 i = 0; i < 24; i += 8) { u64 dgst[8]; diff --git a/OpenCL/m06231.cl b/OpenCL/m06231.cl index 4e2645ed8..918cf699d 100644 --- a/OpenCL/m06231.cl +++ b/OpenCL/m06231.cl @@ -1850,8 +1850,6 @@ __kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf2[14] = 0; salt_buf2[15] = (64 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 w[16]; w[ 0] = swap32 (w0[0]); @@ -1910,7 +1908,7 @@ __kernel void m06231_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[14] = opad[14]; tmps[gid].opad[15] = opad[15]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 16, j += 1) + for (u32 i = 0, j = 1; i < 16; i += 16, j += 1) { salt_buf2[0] = j; @@ -1996,8 +1994,6 @@ __kernel void m06231_loop (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 ipad[16]; ipad[ 0] = tmps[gid].ipad[ 0]; @@ -2036,7 +2032,7 @@ __kernel void m06231_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[14] = tmps[gid].opad[14]; opad[15] = tmps[gid].opad[15]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 16) + for (u32 i = 0; i < 16; i += 16) { u32 dgst[16]; diff --git a/OpenCL/m06232.cl b/OpenCL/m06232.cl index df12a7ab1..099f9aa2b 100644 --- a/OpenCL/m06232.cl +++ b/OpenCL/m06232.cl @@ -1600,8 +1600,6 @@ __kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf2[14] = 0; salt_buf2[15] = (64 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 w[16]; w[ 0] = swap32 (w0[0]); @@ -1660,7 +1658,7 @@ __kernel void m06232_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[14] = opad[14]; tmps[gid].opad[15] = opad[15]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 16, j += 1) + for (u32 i = 0, j = 1; i < 32; i += 16, j += 1) { salt_buf2[0] = j; @@ -1746,8 +1744,6 @@ __kernel void m06232_loop (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 ipad[16]; ipad[ 0] = tmps[gid].ipad[ 0]; @@ -1786,7 +1782,7 @@ __kernel void m06232_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[14] = tmps[gid].opad[14]; opad[15] = tmps[gid].opad[15]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 16) + for (u32 i = 0; i < 32; i += 16) { u32 dgst[16]; diff --git a/OpenCL/m06233.cl b/OpenCL/m06233.cl index aa8306626..1a3e05619 100644 --- a/OpenCL/m06233.cl +++ b/OpenCL/m06233.cl @@ -1600,8 +1600,6 @@ __kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rul salt_buf2[14] = 0; salt_buf2[15] = (64 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 w[16]; w[ 0] = swap32 (w0[0]); @@ -1660,7 +1658,7 @@ __kernel void m06233_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[14] = opad[14]; tmps[gid].opad[15] = opad[15]; - for (u32 i = 0, j = 1; i < (truecrypt_mdlen / 8 / 4); i += 16, j += 1) + for (u32 i = 0, j = 1; i < 48; i += 16, j += 1) { salt_buf2[0] = j; @@ -1746,8 +1744,6 @@ __kernel void m06233_loop (__global pw_t *pws, __global const kernel_rule_t *rul if (gid >= gid_max) return; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 ipad[16]; ipad[ 0] = tmps[gid].ipad[ 0]; @@ -1786,7 +1782,7 @@ __kernel void m06233_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[14] = tmps[gid].opad[14]; opad[15] = tmps[gid].opad[15]; - for (u32 i = 0; i < (truecrypt_mdlen / 8 / 4); i += 16) + for (u32 i = 0; i < 48; i += 16) { u32 dgst[16]; diff --git a/OpenCL/m13751.cl b/OpenCL/m13751.cl index 00657b179..87504ea20 100644 --- a/OpenCL/m13751.cl +++ b/OpenCL/m13751.cl @@ -418,8 +418,6 @@ __kernel void m13751_init (__global pw_t *pws, __global const kernel_rule_t *rul s7[2] = 0; s7[3] = (64 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 ipad[8]; u32 opad[8]; @@ -443,7 +441,7 @@ __kernel void m13751_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[6] = opad[6]; tmps[gid].opad[7] = opad[7]; - for (u32 i = 0, j = 1; i < ((truecrypt_mdlen / 8) / 4); i += 8, j += 1) + for (u32 i = 0, j = 1; i < 16; i += 8, j += 1) { s4[0] = j; @@ -473,8 +471,6 @@ __kernel void m13751_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m13751_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 tc_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -501,7 +497,7 @@ __kernel void m13751_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[6] = tmps[gid].opad[6]; opad[7] = tmps[gid].opad[7]; - for (u32 i = 0; i < ((truecrypt_mdlen / 8) / 4); i += 8) + for (u32 i = 0; i < 16; i += 8) { u32 dgst[8]; diff --git a/OpenCL/m13752.cl b/OpenCL/m13752.cl index 21005b80a..517df9423 100644 --- a/OpenCL/m13752.cl +++ b/OpenCL/m13752.cl @@ -418,8 +418,6 @@ __kernel void m13752_init (__global pw_t *pws, __global const kernel_rule_t *rul s7[2] = 0; s7[3] = (64 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 ipad[8]; u32 opad[8]; @@ -443,7 +441,7 @@ __kernel void m13752_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[6] = opad[6]; tmps[gid].opad[7] = opad[7]; - for (u32 i = 0, j = 1; i < ((truecrypt_mdlen / 8) / 4); i += 8, j += 1) + for (u32 i = 0, j = 1; i < 32; i += 8, j += 1) { s4[0] = j; @@ -473,8 +471,6 @@ __kernel void m13752_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m13752_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 tc_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -501,7 +497,7 @@ __kernel void m13752_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[6] = tmps[gid].opad[6]; opad[7] = tmps[gid].opad[7]; - for (u32 i = 0; i < ((truecrypt_mdlen / 8) / 4); i += 8) + for (u32 i = 0; i < 32; i += 8) { u32 dgst[8]; diff --git a/OpenCL/m13753.cl b/OpenCL/m13753.cl index e7c5acb67..b40643e1c 100644 --- a/OpenCL/m13753.cl +++ b/OpenCL/m13753.cl @@ -418,8 +418,6 @@ __kernel void m13753_init (__global pw_t *pws, __global const kernel_rule_t *rul s7[2] = 0; s7[3] = (64 + 64 + 4) * 8; - const u32 truecrypt_mdlen = salt_bufs[0].truecrypt_mdlen; - u32 ipad[8]; u32 opad[8]; @@ -443,7 +441,7 @@ __kernel void m13753_init (__global pw_t *pws, __global const kernel_rule_t *rul tmps[gid].opad[6] = opad[6]; tmps[gid].opad[7] = opad[7]; - for (u32 i = 0, j = 1; i < ((truecrypt_mdlen / 8) / 4); i += 8, j += 1) + for (u32 i = 0, j = 1; i < 48; i += 8, j += 1) { s4[0] = j; @@ -473,8 +471,6 @@ __kernel void m13753_init (__global pw_t *pws, __global const kernel_rule_t *rul __kernel void m13753_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 tc_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 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; - const u32 gid = get_global_id (0); if (gid >= gid_max) return; @@ -501,7 +497,7 @@ __kernel void m13753_loop (__global pw_t *pws, __global const kernel_rule_t *rul opad[6] = tmps[gid].opad[6]; opad[7] = tmps[gid].opad[7]; - for (u32 i = 0; i < ((truecrypt_mdlen / 8) / 4); i += 8) + for (u32 i = 0; i < 48; i += 8) { u32 dgst[8]; diff --git a/include/types.h b/include/types.h index 96022ff12..fba9e44bf 100644 --- a/include/types.h +++ b/include/types.h @@ -632,7 +632,6 @@ typedef struct salt u32 salt_sign[2]; u32 keccak_mdlen; - u32 truecrypt_mdlen; u32 digests_cnt; u32 digests_done; diff --git a/src/interface.c b/src/interface.c index c69e996c2..07874f51b 100644 --- a/src/interface.c +++ b/src/interface.c @@ -20018,54 +20018,8 @@ u32 hashconfig_get_kernel_loops (hashcat_ctx_t *hashcat_ctx) int hashconfig_general_defaults (hashcat_ctx_t *hashcat_ctx) { - const hashconfig_t *hashconfig = hashcat_ctx->hashconfig; - const hashes_t *hashes = hashcat_ctx->hashes; const user_options_t *user_options = hashcat_ctx->user_options; - salt_t *salts_buf = hashes->salts_buf; - void *esalts_buf = hashes->esalts_buf; - - /** - * special modification not set from parser - */ - - for (u32 salts_pos = 0; salts_pos < hashes->salts_cnt; salts_pos++) - { - switch (hashconfig->hash_mode) - { - case 6211: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 6212: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 6213: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 6221: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 6222: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 6223: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 6231: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 6232: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 6233: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 6241: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 6242: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 6243: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 13711: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 13712: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 13713: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 13721: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 13722: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 13723: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 13731: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 13732: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 13733: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 13741: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 13742: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 13743: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 13751: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 13752: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 13753: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - case 13761: salts_buf[salts_pos].truecrypt_mdlen = 1 * 512; break; - case 13762: salts_buf[salts_pos].truecrypt_mdlen = 2 * 512; break; - case 13763: salts_buf[salts_pos].truecrypt_mdlen = 3 * 512; break; - } - } - char *optional_param1 = NULL; if (user_options->truecrypt_keyfiles) optional_param1 = user_options->truecrypt_keyfiles; @@ -20073,6 +20027,10 @@ int hashconfig_general_defaults (hashcat_ctx_t *hashcat_ctx) if (optional_param1) { + const hashes_t *hashes = hashcat_ctx->hashes; + + void *esalts_buf = hashes->esalts_buf; + char *tcvc_keyfiles = (char *) optional_param1; u32 *keyfile_buf = ((tc_t *) esalts_buf)->keyfile_buf;