github/hashcat
1
Fork 0
mirror of https://github.com/hashcat/hashcat synced 2024-12-16 08:15:54 +01:00
Code Releases Activity

Get rid of truecrypt_mdlen

Browse Source
This commit is contained in:
jsteube 2016-11-30 11:41:31 +01:00
parent 3c40b88eff
commit cad3b3e10b
15 changed files with 28 additions and 120 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
OpenCL/inc_types.cl
View File

@ -677,7 +677,6 @@ typedef struct
u32 salt_sign[2];
u32 keccak_mdlen;
u32 truecrypt_mdlen;
u32 digests_cnt;
u32 digests_done;

8
OpenCL/m06211.cl
View File

@ -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];

8
OpenCL/m06212.cl
View File

@ -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];

8
OpenCL/m06213.cl
View File

@ -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];

8
OpenCL/m06221.cl
View File

@ -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];

8
OpenCL/m06222.cl
View File

@ -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];

8
OpenCL/m06223.cl
View File

@ -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];

8
OpenCL/m06231.cl
View File

@ -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];

8
OpenCL/m06232.cl
View File

@ -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];

8
OpenCL/m06233.cl
View File

@ -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];

8
OpenCL/m13751.cl
View File

@ -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];

8
OpenCL/m13752.cl
View File

@ -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];

8
OpenCL/m13753.cl
View File

@ -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];

1
include/types.h
View File

@ -632,7 +632,6 @@ typedef struct salt
u32 salt_sign[2];
u32 keccak_mdlen;
u32 truecrypt_mdlen;
u32 digests_cnt;
u32 digests_done;

50
src/interface.c
View File

@ -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;