arno/hashcat
1
0
Fork 0
mirror of https://github.com/hashcat/hashcat.git synced 2024-11-22 16:18:09 +00:00
Code Issues Releases Wiki Activity

Remove unused truncate_block_xxx_xx() functions and update kernels to use the _S function

Browse Source
This commit is contained in:
jsteube 2017-08-24 20:07:43 +02:00
parent 51dc1c7db3
commit bed7e8f466
17 changed files with 1763 additions and 1873 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3554
OpenCL/inc_common.cl
View File

File diff suppressed because it is too large Load Diff

2
OpenCL/m00500-optimized.cl
View File

@ -776,7 +776,7 @@ __kernel void m00500_init (__global pw_t *pws, __constant const kernel_rule_t *r
block_len += salt_len;
truncate_block_4x4_le (digest, pw_len);
truncate_block_4x4_le_S (digest, pw_len);
memcat16 (block0, block1, block2, block3, block_len, digest);

2
OpenCL/m00500.cl
View File

@ -93,7 +93,7 @@ __kernel void m00500_init (__global pw_t *pws, __constant const kernel_rule_t *r
md5_update (&md5_ctx, final, 16);
}
truncate_block_4x4_le (final, pl);
truncate_block_4x4_le_S (final, pl);
md5_update (&md5_ctx, final, pl);

2
OpenCL/m01600-optimized.cl
View File

@ -791,7 +791,7 @@ __kernel void m01600_init (__global pw_t *pws, __constant const kernel_rule_t *r
block_len += salt_len;
truncate_block_4x4_le (digest, pw_len);
truncate_block_4x4_le_S (digest, pw_len);
memcat16 (block0, block1, block2, block3, block_len, digest);

2
OpenCL/m01600.cl
View File

@ -95,7 +95,7 @@ __kernel void m01600_init (__global pw_t *pws, __constant const kernel_rule_t *r
md5_update (&md5_ctx, final, 16);
}
truncate_block_4x4_le (final, pl);
truncate_block_4x4_le_S (final, pl);
md5_update (&md5_ctx, final, pl);

6
OpenCL/m01800.cl
View File

@ -112,7 +112,7 @@ __kernel void m01800_init (__global pw_t *pws, __constant const kernel_rule_t *r
#endif
for (int i = 0; i < 16; i++) t_final[i] = final[i];
truncate_block_16x4_be (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
truncate_block_16x4_be_S (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
sha512_update (&ctx, t_final, pl);
@ -199,7 +199,7 @@ __kernel void m01800_init (__global pw_t *pws, __constant const kernel_rule_t *r
p_final[idx + 15] = final[15];
}
truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
truncate_block_16x4_be_S (final + 0, final + 4, final + 8, final + 12, pl);
p_final[idx + 0] = final[ 0];
p_final[idx + 1] = final[ 1];
@ -273,7 +273,7 @@ __kernel void m01800_init (__global pw_t *pws, __constant const kernel_rule_t *r
s_final[idx + 15] = final[15];
}
truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
truncate_block_16x4_be_S (final + 0, final + 4, final + 8, final + 12, pl);
s_final[idx + 0] = final[ 0];
s_final[idx + 1] = final[ 1];

2
OpenCL/m06300-optimized.cl
View File

@ -685,7 +685,7 @@ __kernel void m06300_init (__global pw_t *pws, __constant const kernel_rule_t *r
block_len += salt_len;
truncate_block_4x4_le (digest, pw_len);
truncate_block_4x4_le_S (digest, pw_len);
memcat16 (block0, block1, block2, block3, block_len, digest);

2
OpenCL/m06300.cl
View File

@ -85,7 +85,7 @@ __kernel void m06300_init (__global pw_t *pws, __constant const kernel_rule_t *r
md5_update (&md5_ctx, final, 16);
}
truncate_block_4x4_le (final, pl);
truncate_block_4x4_le_S (final, pl);
md5_update (&md5_ctx, final, pl);

6
OpenCL/m07400-optimized.cl
View File

@ -780,7 +780,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
alt_result_tmp[6] = 0;
alt_result_tmp[7] = 0;
truncate_block_4x4_le (alt_result_tmp, pw_len);
truncate_block_4x4_le_S (alt_result_tmp, pw_len);
/* Add the key string. */
@ -890,7 +890,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
bswap8 (p_bytes);
truncate_block_4x4_le (p_bytes, pw_len);
truncate_block_4x4_le_S (p_bytes, pw_len);
tmps[gid].p_bytes[0] = p_bytes[0];
tmps[gid].p_bytes[1] = p_bytes[1];
@ -933,7 +933,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
bswap8 (s_bytes);
truncate_block_4x4_le (s_bytes, salt_len);
truncate_block_4x4_le_S (s_bytes, salt_len);
tmps[gid].s_bytes[0] = s_bytes[0];
tmps[gid].s_bytes[1] = s_bytes[1];

6
OpenCL/m07400.cl
View File

@ -104,7 +104,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
#endif
for (int i = 0; i < 8; i++) t_final[i] = final[i];
truncate_block_16x4_be (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
truncate_block_16x4_be_S (t_final + 0, t_final + 4, t_final + 8, t_final + 12, pl);
sha256_update (&ctx, t_final, pl);
@ -175,7 +175,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
p_final[idx + 7] = final[7];
}
truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
truncate_block_16x4_be_S (final + 0, final + 4, final + 8, final + 12, pl);
p_final[idx + 0] = final[0];
p_final[idx + 1] = final[1];
@ -233,7 +233,7 @@ __kernel void m07400_init (__global pw_t *pws, __constant const kernel_rule_t *r
s_final[idx + 7] = final[7];
}
truncate_block_16x4_be (final + 0, final + 4, final + 8, final + 12, pl);
truncate_block_16x4_be_S (final + 0, final + 4, final + 8, final + 12, pl);
s_final[idx + 0] = final[0];
s_final[idx + 1] = final[1];

8
OpenCL/m13100_a0-optimized.cl
View File

@ -321,7 +321,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block_4x4_le (w0, edata2_left & 0xf);
truncate_block_4x4_le_S (w0, edata2_left & 0xf);
append_0x80_1x4 (w0, edata2_left & 0xf);
@ -335,7 +335,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block_4x4_le (w1, edata2_left & 0xf);
truncate_block_4x4_le_S (w1, edata2_left & 0xf);
append_0x80_1x4 (w1, edata2_left & 0xf);
@ -350,7 +350,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block_4x4_le (w2, edata2_left & 0xf);
truncate_block_4x4_le_S (w2, edata2_left & 0xf);
append_0x80_1x4 (w2, edata2_left & 0xf);
@ -366,7 +366,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block_4x4_le (w3, edata2_left & 0xf);
truncate_block_4x4_le_S (w3, edata2_left & 0xf);
append_0x80_1x4 (w3, edata2_left & 0xf);

8
OpenCL/m13100_a0.cl
View File

@ -234,14 +234,14 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block_4x4_le (w0, edata2_left & 0xf);
truncate_block_4x4_le_S (w0, edata2_left & 0xf);
}
else if (edata2_left < 32)
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block_4x4_le (w1, edata2_left & 0xf);
truncate_block_4x4_le_S (w1, edata2_left & 0xf);
}
else if (edata2_left < 48)
{
@ -249,7 +249,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block_4x4_le (w2, edata2_left & 0xf);
truncate_block_4x4_le_S (w2, edata2_left & 0xf);
}
else
{
@ -258,7 +258,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block_4x4_le (w3, edata2_left & 0xf);
truncate_block_4x4_le_S (w3, edata2_left & 0xf);
}
md5_hmac_update_64 (&ctx, w0, w1, w2, w3, edata2_left);

8
OpenCL/m13100_a1-optimized.cl
View File

@ -319,7 +319,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block_4x4_le (w0, edata2_left & 0xf);
truncate_block_4x4_le_S (w0, edata2_left & 0xf);
append_0x80_1x4 (w0, edata2_left & 0xf);
@ -333,7 +333,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block_4x4_le (w1, edata2_left & 0xf);
truncate_block_4x4_le_S (w1, edata2_left & 0xf);
append_0x80_1x4 (w1, edata2_left & 0xf);
@ -348,7 +348,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block_4x4_le (w2, edata2_left & 0xf);
truncate_block_4x4_le_S (w2, edata2_left & 0xf);
append_0x80_1x4 (w2, edata2_left & 0xf);
@ -364,7 +364,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block_4x4_le (w3, edata2_left & 0xf);
truncate_block_4x4_le_S (w3, edata2_left & 0xf);
append_0x80_1x4 (w3, edata2_left & 0xf);

8
OpenCL/m13100_a1.cl
View File

@ -232,14 +232,14 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block_4x4_le (w0, edata2_left & 0xf);
truncate_block_4x4_le_S (w0, edata2_left & 0xf);
}
else if (edata2_left < 32)
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block_4x4_le (w1, edata2_left & 0xf);
truncate_block_4x4_le_S (w1, edata2_left & 0xf);
}
else if (edata2_left < 48)
{
@ -247,7 +247,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block_4x4_le (w2, edata2_left & 0xf);
truncate_block_4x4_le_S (w2, edata2_left & 0xf);
}
else
{
@ -256,7 +256,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block_4x4_le (w3, edata2_left & 0xf);
truncate_block_4x4_le_S (w3, edata2_left & 0xf);
}
md5_hmac_update_64 (&ctx, w0, w1, w2, w3, edata2_left);

8
OpenCL/m13100_a3-optimized.cl
View File

@ -319,7 +319,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block_4x4_le (w0, edata2_left & 0xf);
truncate_block_4x4_le_S (w0, edata2_left & 0xf);
append_0x80_1x4 (w0, edata2_left & 0xf);
@ -333,7 +333,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block_4x4_le (w1, edata2_left & 0xf);
truncate_block_4x4_le_S (w1, edata2_left & 0xf);
append_0x80_1x4 (w1, edata2_left & 0xf);
@ -348,7 +348,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block_4x4_le (w2, edata2_left & 0xf);
truncate_block_4x4_le_S (w2, edata2_left & 0xf);
append_0x80_1x4 (w2, edata2_left & 0xf);
@ -364,7 +364,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global u32 *edat
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block_4x4_le (w3, edata2_left & 0xf);
truncate_block_4x4_le_S (w3, edata2_left & 0xf);
append_0x80_1x4 (w3, edata2_left & 0xf);

8
OpenCL/m13100_a3.cl
View File

@ -232,14 +232,14 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
truncate_block_4x4_le (w0, edata2_left & 0xf);
truncate_block_4x4_le_S (w0, edata2_left & 0xf);
}
else if (edata2_left < 32)
{
j = rc4_next_16 (rc4_key, i, j, edata2, w0); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
truncate_block_4x4_le (w1, edata2_left & 0xf);
truncate_block_4x4_le_S (w1, edata2_left & 0xf);
}
else if (edata2_left < 48)
{
@ -247,7 +247,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
j = rc4_next_16 (rc4_key, i, j, edata2, w1); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
truncate_block_4x4_le (w2, edata2_left & 0xf);
truncate_block_4x4_le_S (w2, edata2_left & 0xf);
}
else
{
@ -256,7 +256,7 @@ int decrypt_and_check (__local RC4_KEY *rc4_key, u32 data[4], __global const u32
j = rc4_next_16 (rc4_key, i, j, edata2, w2); i += 16; edata2 += 4;
j = rc4_next_16 (rc4_key, i, j, edata2, w3); i += 16; edata2 += 4;
truncate_block_4x4_le (w3, edata2_left & 0xf);
truncate_block_4x4_le_S (w3, edata2_left & 0xf);
}
md5_hmac_update_64 (&ctx, w0, w1, w2, w3, edata2_left);

4
OpenCL/m13400.cl
View File

@ -524,7 +524,7 @@ __kernel void m13400_comp (__global pw_t *pws, __constant const kernel_rule_t *r
// we need to clear the buffer of the padding data
truncate_block_4x4_be (out, 16 - pad_byte);
truncate_block_4x4_be_S (out, 16 - pad_byte);
u32 w0[4] = { 0 };
u32 w1[4] = { 0 };
@ -613,7 +613,7 @@ __kernel void m13400_comp (__global pw_t *pws, __constant const kernel_rule_t *r
// we need to clear the buffer of the padding data
truncate_block_4x4_be (out, 16 - pad_byte);
truncate_block_4x4_be_S (out, 16 - pad_byte);
u32 w0[4] = { 0 };
u32 w1[4] = { 0 };