From: jsteube Date: Tue, 3 May 2016 08:57:23 +0000 (+0200) Subject: Fix -m 1500 multihash and benchmark X-Git-Tag: v3.00-beta~47 X-Git-Url: https://www.flypig.org.uk/git/?p=hashcat.git;a=commitdiff_plain;h=dbe2d966188bbfffd0f397c3e3ea449104d3bfbb Fix -m 1500 multihash and benchmark 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 --- diff --git a/OpenCL/m01500_a0.cl b/OpenCL/m01500_a0.cl index 3a0e663..6bf6f34 100644 --- a/OpenCL/m01500_a0.cl +++ b/OpenCL/m01500_a0.cl @@ -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); - u32x c = 0; - u32x d = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 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); - u32x c = 0; - u32x d = 0; + u32x z = 0; - COMPARE_S_SIMD (iv[0], iv[1], c, d); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m01500_a1.cl b/OpenCL/m01500_a1.cl index b374425..4d7a547 100644 --- a/OpenCL/m01500_a1.cl +++ b/OpenCL/m01500_a1.cl @@ -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); - u32x c = 0; - u32x d = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 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); - u32x c = 0; - u32x d = 0; + u32x z = 0; - COMPARE_S_SIMD (iv[0], iv[1], c, d); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m03000_a0.cl b/OpenCL/m03000_a0.cl index 5be3695..d2aa49d 100644 --- a/OpenCL/m03000_a0.cl +++ b/OpenCL/m03000_a0.cl @@ -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); - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 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); - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 0; - COMPARE_S_SIMD (r0, r1, r2, r3); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m03000_a1.cl b/OpenCL/m03000_a1.cl index 8b7283c..7079a7a 100644 --- a/OpenCL/m03000_a1.cl +++ b/OpenCL/m03000_a1.cl @@ -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); - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 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); - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 0; - COMPARE_S_SIMD (r0, r1, r2, r3); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m03100_a0.cl b/OpenCL/m03100_a0.cl index 85e91f7..9220eb7 100644 --- a/OpenCL/m03100_a0.cl +++ b/OpenCL/m03100_a0.cl @@ -703,12 +703,9 @@ __kernel void m03100_m04 (__global pw_t *pws, __global kernel_rule_t * rules_bu * cmp */ - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** @@ -941,12 +938,9 @@ __kernel void m03100_s04 (__global pw_t *pws, __global kernel_rule_t * rules_bu * cmp */ - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 0; - COMPARE_S_SIMD (r0, r1, r2, r3); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m03100_a1.cl b/OpenCL/m03100_a1.cl index a0a09c7..5f3207a 100644 --- a/OpenCL/m03100_a1.cl +++ b/OpenCL/m03100_a1.cl @@ -760,12 +760,9 @@ __kernel void m03100_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, * cmp */ - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** @@ -1058,12 +1055,9 @@ __kernel void m03100_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, * cmp */ - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 0; - COMPARE_S_SIMD (r0, r1, r2, r3); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m03100_a3.cl b/OpenCL/m03100_a3.cl index c7a7b68..63a7eea 100644 --- a/OpenCL/m03100_a3.cl +++ b/OpenCL/m03100_a3.cl @@ -681,12 +681,9 @@ void m03100m (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], * cmp */ - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** @@ -892,12 +889,9 @@ void m03100s (__local u32 (*s_SPtrans)[64], __local u32 (*s_skb)[64], u32 w[16], * cmp */ - const u32x r0 = iv[0]; - const u32x r1 = iv[1]; - const u32x r2 = 0; - const u32x r3 = 0; + u32x z = 0; - COMPARE_S_SIMD (r0, r1, r2, r3); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m07700_a0.cl b/OpenCL/m07700_a0.cl index ebe418b..2755132 100644 --- a/OpenCL/m07700_a0.cl +++ b/OpenCL/m07700_a0.cl @@ -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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m07700_a1.cl b/OpenCL/m07700_a1.cl index 28c6827..8eb9054 100644 --- a/OpenCL/m07700_a1.cl +++ b/OpenCL/m07700_a1.cl @@ -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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m07700_a3.cl b/OpenCL/m07700_a3.cl index 12470ff..dac0379 100644 --- a/OpenCL/m07700_a3.cl +++ b/OpenCL/m07700_a3.cl @@ -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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0 }; /** diff --git a/OpenCL/m08500_a0.cl b/OpenCL/m08500_a0.cl index 71b0e00..8fccfd6 100644 --- a/OpenCL/m08500_a0.cl +++ b/OpenCL/m08500_a0.cl @@ -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); - u32x iv2 = 0; - u32x iv3 = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 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); - u32x iv2 = 0; - u32x iv3 = 0; + u32x z = 0; - COMPARE_S_SIMD (iv[0], iv[1], iv2, iv3); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m08500_a1.cl b/OpenCL/m08500_a1.cl index bd7e44c..63ee260 100644 --- a/OpenCL/m08500_a1.cl +++ b/OpenCL/m08500_a1.cl @@ -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); - u32x iv2 = 0; - u32x iv3 = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 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); - u32x iv2 = 0; - u32x iv3 = 0; + u32x z = 0; - COMPARE_S_SIMD (iv[0], iv[1], iv2, iv3); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m08500_a3.cl b/OpenCL/m08500_a3.cl index 686bf93..933261a 100644 --- a/OpenCL/m08500_a3.cl +++ b/OpenCL/m08500_a3.cl @@ -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); - u32x iv2 = 0; - u32x iv3 = 0; + u32x z = 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_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 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); - u32x iv2 = 0; - u32x iv3 = 0; + u32x z = 0; - COMPARE_S_SIMD (iv[0], iv[1], iv2, iv3); + COMPARE_S_SIMD (iv[0], iv[1], z, z); } } diff --git a/OpenCL/m11500_a0.cl b/OpenCL/m11500_a0.cl index 4feacf9..3bcc15b 100644 --- a/OpenCL/m11500_a0.cl +++ b/OpenCL/m11500_a0.cl @@ -209,11 +209,10 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, w[15] = 0; 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] = { digests_buf[digests_offset].digest_buf[DGST_R0], - digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0, + 0 }; /** @@ -310,11 +309,10 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, w[15] = 0; 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); } } diff --git a/OpenCL/m11500_a1.cl b/OpenCL/m11500_a1.cl index 3182074..3dd6c46 100644 --- a/OpenCL/m11500_a1.cl +++ b/OpenCL/m11500_a1.cl @@ -265,11 +265,10 @@ __kernel void m11500_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, w[15] = w3[3]; 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] = { digests_buf[digests_offset].digest_buf[DGST_R0], - digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0, + 0 }; /** @@ -426,11 +425,10 @@ __kernel void m11500_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf, w[15] = w3[3]; 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); } } diff --git a/OpenCL/m11500_a3.cl b/OpenCL/m11500_a3.cl index d1d483e..c8e9986 100644 --- a/OpenCL/m11500_a3.cl +++ b/OpenCL/m11500_a3.cl @@ -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]; 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] = { digests_buf[digests_offset].digest_buf[DGST_R0], - digests_buf[digests_offset].digest_buf[DGST_R1], - digests_buf[digests_offset].digest_buf[DGST_R2], - digests_buf[digests_offset].digest_buf[DGST_R3] + 0, + 0, + 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]; 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); } }