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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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 (file)
--- 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);
   }
 }