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

Fix -m 1500 multihash and benchmark

Browse Source 
Fix -m  3000 multihash and benchmark
Fix -m  3100 multihash and benchmark
Fix -m  7700 multihash and benchmark
Fix -m  8500 multihash and benchmark
Fix -m 11500 multihash and benchmark
This commit is contained in:
jsteube 2016-05-03 10:57:23 +02:00
parent c7b67376a8
commit dbe2d96618
16 changed files with 93 additions and 139 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
OpenCL/m01500_a0.cl
View File

@ -580,10 +580,9 @@ __kernel void m01500_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
_des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
u32x c = 0; u32x z = 0;
u32x d = 0;
COMPARE_M_SIMD (iv[0], iv[1], c, d); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -669,8 +668,8 @@ __kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -700,10 +699,9 @@ __kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
_des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
u32x c = 0; u32x z = 0;
u32x d = 0;
COMPARE_S_SIMD (iv[0], iv[1], c, d); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

14
OpenCL/m01500_a1.cl
View File

@ -641,10 +641,9 @@ __kernel void m01500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
_des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
u32x c = 0; u32x z = 0;
u32x d = 0;
COMPARE_M_SIMD (iv[0], iv[1], c, d); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -730,8 +729,8 @@ __kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -825,10 +824,9 @@ __kernel void m01500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
_des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, mask, Kc, Kd, s_SPtrans);
u32x c = 0; u32x z = 0;
u32x d = 0;
COMPARE_S_SIMD (iv[0], iv[1], c, d); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

18
OpenCL/m03000_a0.cl
View File

@ -592,12 +592,9 @@ __kernel void m03000_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_M_SIMD (r0, r1, r2, r3); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -677,8 +674,8 @@ __kernel void m03000_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -715,12 +712,9 @@ __kernel void m03000_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_S_SIMD (r0, r1, r2, r3); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

18
OpenCL/m03000_a1.cl
View File

@ -653,12 +653,9 @@ __kernel void m03000_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_M_SIMD (r0, r1, r2, r3); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -738,8 +735,8 @@ __kernel void m03000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -840,12 +837,9 @@ __kernel void m03000_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_S_SIMD (r0, r1, r2, r3); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

18
OpenCL/m03100_a0.cl
View File

@ -703,12 +703,9 @@ __kernel void m03100_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
* cmp * cmp
*/ */
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_M_SIMD (r0, r1, r2, r3); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -806,8 +803,8 @@ __kernel void m03100_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -941,12 +938,9 @@ __kernel void m03100_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
* cmp * cmp
*/ */
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_S_SIMD (r0, r1, r2, r3); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

18
OpenCL/m03100_a1.cl
View File

@ -760,12 +760,9 @@ __kernel void m03100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
* cmp * cmp
*/ */
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_M_SIMD (r0, r1, r2, r3); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -863,8 +860,8 @@ __kernel void m03100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -1058,12 +1055,9 @@ __kernel void m03100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
* cmp * cmp
*/ */
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_S_SIMD (r0, r1, r2, r3); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

18
OpenCL/m03100_a3.cl
View File

@ -681,12 +681,9 @@ void m03100m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16],
* cmp * cmp
*/ */
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_M_SIMD (r0, r1, r2, r3); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -783,8 +780,8 @@ void m03100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16],
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -892,12 +889,9 @@ void m03100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16],
* cmp * cmp
*/ */
const u32x r0 = iv[0]; u32x z = 0;
const u32x r1 = iv[1];
const u32x r2 = 0;
const u32x r3 = 0;
COMPARE_S_SIMD (r0, r1, r2, r3); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

4
OpenCL/m07700_a0.cl
View File

@ -577,8 +577,8 @@ __kernel void m07700_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**

4
OpenCL/m07700_a1.cl
View File

@ -618,8 +618,8 @@ __kernel void m07700_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**

4
OpenCL/m07700_a3.cl
View File

@ -543,8 +543,8 @@ void m07700s (u32 w0[4], u32 w1[4], u32 w2[4], u32 w3[4], const u32 pw_len, __gl
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**

14
OpenCL/m08500_a0.cl
View File

@ -635,10 +635,9 @@ __kernel void m08500_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x iv2 = 0; u32x z = 0;
u32x iv3 = 0;
COMPARE_M_SIMD (iv[0], iv[1], iv2, iv3); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -727,8 +726,8 @@ __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -769,10 +768,9 @@ __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x iv2 = 0; u32x z = 0;
u32x iv3 = 0;
COMPARE_S_SIMD (iv[0], iv[1], iv2, iv3); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

14
OpenCL/m08500_a1.cl
View File

@ -675,10 +675,9 @@ __kernel void m08500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x iv2 = 0; u32x z = 0;
u32x iv3 = 0;
COMPARE_M_SIMD (iv[0], iv[1], iv2, iv3); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -767,8 +766,8 @@ __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -852,10 +851,9 @@ __kernel void m08500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x iv2 = 0; u32x z = 0;
u32x iv3 = 0;
COMPARE_S_SIMD (iv[0], iv[1], iv2, iv3); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

14
OpenCL/m08500_a3.cl
View File

@ -583,10 +583,9 @@ void m08500m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16],
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x iv2 = 0; u32x z = 0;
u32x iv3 = 0;
COMPARE_M_SIMD (iv[0], iv[1], iv2, iv3); COMPARE_M_SIMD (iv[0], iv[1], z, z);
} }
} }
@ -616,8 +615,8 @@ void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16],
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], digests_buf[digests_offset].digest_buf[DGST_R1],
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -659,10 +658,9 @@ void m08500s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16],
_des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans); _des_crypt_encrypt (iv, data, Kc, Kd, s_SPtrans);
u32x iv2 = 0; u32x z = 0;
u32x iv3 = 0;
COMPARE_S_SIMD (iv[0], iv[1], iv2, iv3); COMPARE_S_SIMD (iv[0], iv[1], z, z);
} }
} }

20
OpenCL/m11500_a0.cl
View File

@ -209,11 +209,10 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
w[15] = 0; w[15] = 0;
u32x a = crc32 (w, out_len, iv); u32x a = crc32 (w, out_len, iv);
u32x b = 0;
u32x c = 0;
u32x d = 0;
COMPARE_M_SIMD (a, b, c, d); u32x z = 0;
COMPARE_M_SIMD (a, z, z, z);
} }
} }
@ -268,9 +267,9 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
const u32 search[4] = const u32 search[4] =
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], 0,
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -310,11 +309,10 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
w[15] = 0; w[15] = 0;
u32x a = crc32 (w, out_len, iv); u32x a = crc32 (w, out_len, iv);
u32x b = 0;
u32x c = 0;
u32x d = 0;
COMPARE_S_SIMD (a, b, c, d); u32x z = 0;
COMPARE_S_SIMD (a, z, z, z);
} }
} }

20
OpenCL/m11500_a1.cl
View File

@ -265,11 +265,10 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
w[15] = w3[3]; w[15] = w3[3];
u32x a = crc32 (w, pw_len, iv); u32x a = crc32 (w, pw_len, iv);
u32x b = 0;
u32x c = 0;
u32x d = 0;
COMPARE_M_SIMD (a, b, c, d); u32x z = 0;
COMPARE_M_SIMD (a, z, z, z);
} }
} }
@ -324,9 +323,9 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
const u32 search[4] = const u32 search[4] =
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], 0,
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -426,11 +425,10 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
w[15] = w3[3]; w[15] = w3[3];
u32x a = crc32 (w, pw_len, iv); u32x a = crc32 (w, pw_len, iv);
u32x b = 0;
u32x c = 0;
u32x d = 0;
COMPARE_S_SIMD (a, b, c, d); u32x z = 0;
COMPARE_S_SIMD (a, z, z, z);
} }
} }

20
OpenCL/m11500_a3.cl
View File

@ -185,11 +185,10 @@ void m11500m (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_r
w_t[15] = w[15]; w_t[15] = w[15];
u32x a = crc32 (w_t, pw_len, iv); u32x a = crc32 (w_t, pw_len, iv);
u32x b = 0;
u32x c = 0;
u32x d = 0;
COMPARE_M_SIMD (a, b, c, d); u32x z = 0;
COMPARE_M_SIMD (a, z, z, z);
} }
} }
@ -215,9 +214,9 @@ void m11500s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_r
const u32 search[4] = const u32 search[4] =
{ {
digests_buf[digests_offset].digest_buf[DGST_R0], digests_buf[digests_offset].digest_buf[DGST_R0],
digests_buf[digests_offset].digest_buf[DGST_R1], 0,
digests_buf[digests_offset].digest_buf[DGST_R2], 0,
digests_buf[digests_offset].digest_buf[DGST_R3] 0
}; };
/** /**
@ -256,11 +255,10 @@ void m11500s (u32 w[16], const u32 pw_len, __global pw_t *pws, __global kernel_r
w_t[15] = w[15]; w_t[15] = w[15];
u32x a = crc32 (w_t, pw_len, iv); u32x a = crc32 (w_t, pw_len, iv);
u32x b = 0;
u32x c = 0;
u32x d = 0;
COMPARE_S_SIMD (a, b, c, d); u32x z = 0;
COMPARE_S_SIMD (a, z, z, z);
} }
} }