Fix some bugs related to -a0, all caused by the same pitfall

diff --git a/OpenCL/m04700_a0.cl b/OpenCL/m04700_a0.cl
index 1854e19..8a9b5a2 100644 (file)
--- a/OpenCL/m04700_a0.cl
+++ b/OpenCL/m04700_a0.cl
@@ -113,13 +113,15 @@ __kernel void m04700_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     w3[0] = 0;
     w3[1] = 0;
-    w3[2] = pw_len * 8;
+    w3[2] = 0;
     w3[3] = 0;
 
     const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
 
     append_0x80_2x4 (w0, w1, out_len);
 
+    w3[2] = out_len * 8;
+
     /**
      * md5
      */
@@ -451,13 +453,15 @@ __kernel void m04700_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     w3[0] = 0;
     w3[1] = 0;
-    w3[2] = pw_len * 8;
+    w3[2] = 0;
     w3[3] = 0;
 
     const u32 out_len = apply_rules (rules_buf[il_pos].cmds, w0, w1, pw_len);
 
     append_0x80_2x4 (w0, w1, out_len);
 
+    w3[2] = out_len * 8;
+
     /**
      * md5
      */

diff --git a/OpenCL/m06900_a0.cl b/OpenCL/m06900_a0.cl
index a0f13d2..1b15632 100644 (file)
--- a/OpenCL/m06900_a0.cl
+++ b/OpenCL/m06900_a0.cl
@@ -846,7 +846,7 @@ __kernel void m06900_m04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
 
     u32 tmp[8];
 
-    if (pw_len > 0)
+    if (out_len > 0)
     {
       PASS0 (state, tmp, state_m, data_m, s_tables);
       PASS2 (state, tmp, state_m, data_m, s_tables);
@@ -1106,7 +1106,7 @@ __kernel void m06900_s04 (__global pw_t *pws, __global kernel_rule_t *  rules_bu
 
     u32 tmp[8];
 
-    if (pw_len > 0)
+    if (out_len > 0)
     {
       PASS0 (state, tmp, state_m, data_m, s_tables);
       PASS2 (state, tmp, state_m, data_m, s_tables);

diff --git a/OpenCL/m07600_a0.cl b/OpenCL/m07600_a0.cl
index de87dc6..5862cda 100644 (file)
--- a/OpenCL/m07600_a0.cl
+++ b/OpenCL/m07600_a0.cl
@@ -174,7 +174,7 @@ __kernel void m07600_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     u32 wc_t = swap32 (w3[0]);
     u32 wd_t = swap32 (w3[1]);
     u32 we_t = 0;
-    u32 wf_t = pw_len * 8;
+    u32 wf_t = out_len * 8;
 
     u32 a = SHA1M_A;
     u32 b = SHA1M_B;
@@ -769,7 +769,7 @@ __kernel void m07600_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     u32 wc_t = swap32 (w3[0]);
     u32 wd_t = swap32 (w3[1]);
     u32 we_t = 0;
-    u32 wf_t = pw_len * 8;
+    u32 wf_t = out_len * 8;
 
     u32 a = SHA1M_A;
     u32 b = SHA1M_B;

diff --git a/OpenCL/m08300_a0.cl b/OpenCL/m08300_a0.cl
index 2872e9f..3346c64 100644 (file)
--- a/OpenCL/m08300_a0.cl
+++ b/OpenCL/m08300_a0.cl
@@ -280,7 +280,7 @@ __kernel void m08300_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     switch_buffer_by_offset_le (w0_t, w1_t, w2_t, w3_t, 1);
 
-    w0_t[0] |= pw_len & 0xff;
+    w0_t[0] |= out_len & 0xff;
 
     /**
      * salt
@@ -588,7 +588,7 @@ __kernel void m08300_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
 
     switch_buffer_by_offset_le (w0_t, w1_t, w2_t, w3_t, 1);
 
-    w0_t[0] |= pw_len & 0xff;
+    w0_t[0] |= out_len & 0xff;
 
     /**
      * salt

diff --git a/OpenCL/m08400_a0.cl b/OpenCL/m08400_a0.cl
index 348c994..2c29b41 100644 (file)
--- a/OpenCL/m08400_a0.cl
+++ b/OpenCL/m08400_a0.cl
@@ -300,7 +300,7 @@ __kernel void m08400_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     w3_t[0] = swap32 (w3[0]);
     w3_t[1] = swap32 (w3[1]);
     w3_t[2] = 0;
-    w3_t[3] = pw_len * 8;
+    w3_t[3] = out_len * 8;
 
     u32 digest[5];
 
@@ -607,7 +607,7 @@ __kernel void m08400_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
     w3_t[0] = swap32 (w3[0]);
     w3_t[1] = swap32 (w3[1]);
     w3_t[2] = 0;
-    w3_t[3] = pw_len * 8;
+    w3_t[3] = out_len * 8;
 
     u32 digest[5];
 

diff --git a/OpenCL/m08600_a0.cl b/OpenCL/m08600_a0.cl
index 1c688fa..bae3f4c 100644 (file)
--- a/OpenCL/m08600_a0.cl
+++ b/OpenCL/m08600_a0.cl
@@ -345,24 +345,24 @@ __kernel void m08600_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
      * padding
      */
 
-    if (pw_len < 16)
+    if (out_len < 16)
     {
-      pad (&w[ 0], pw_len & 0xf);
+      pad (&w[ 0], out_len & 0xf);
     }
-    else if (pw_len < 32)
+    else if (out_len < 32)
     {
-      pad (&w[ 4], pw_len & 0xf);
+      pad (&w[ 4], out_len & 0xf);
     }
-    else if (pw_len < 48)
+    else if (out_len < 48)
     {
-      pad (&w[ 8], pw_len & 0xf);
+      pad (&w[ 8], out_len & 0xf);
     }
-    else if (pw_len < 64)
+    else if (out_len < 64)
     {
-      pad (&w[12], pw_len & 0xf);
+      pad (&w[12], out_len & 0xf);
     }
 
-    domino_big_md (w, pw_len, state, s_lotus_magic_table);
+    domino_big_md (w, out_len, state, s_lotus_magic_table);
 
     const u32 r0 = state[0];
     const u32 r1 = state[1];
@@ -504,24 +504,24 @@ __kernel void m08600_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
      * padding
      */
 
-    if (pw_len < 16)
+    if (out_len < 16)
     {
-      pad (&w[ 0], pw_len & 0xf);
+      pad (&w[ 0], out_len & 0xf);
     }
-    else if (pw_len < 32)
+    else if (out_len < 32)
     {
-      pad (&w[ 4], pw_len & 0xf);
+      pad (&w[ 4], out_len & 0xf);
     }
-    else if (pw_len < 48)
+    else if (out_len < 48)
     {
-      pad (&w[ 8], pw_len & 0xf);
+      pad (&w[ 8], out_len & 0xf);
     }
-    else if (pw_len < 64)
+    else if (out_len < 64)
     {
-      pad (&w[12], pw_len & 0xf);
+      pad (&w[12], out_len & 0xf);
     }
 
-    domino_big_md (w, pw_len, state, s_lotus_magic_table);
+    domino_big_md (w, out_len, state, s_lotus_magic_table);
 
     const u32 r0 = state[0];
     const u32 r1 = state[1];

diff --git a/OpenCL/m08700_a0.cl b/OpenCL/m08700_a0.cl
index 6376422..91785fc 100644 (file)
--- a/OpenCL/m08700_a0.cl
+++ b/OpenCL/m08700_a0.cl
@@ -401,24 +401,24 @@ __kernel void m08700_m04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
      * padding
      */
 
-    if (pw_len < 16)
+    if (out_len < 16)
     {
-      pad (&w[ 0], pw_len & 0xf);
+      pad (&w[ 0], out_len & 0xf);
     }
-    else if (pw_len < 32)
+    else if (out_len < 32)
     {
-      pad (&w[ 4], pw_len & 0xf);
+      pad (&w[ 4], out_len & 0xf);
     }
-    else if (pw_len < 48)
+    else if (out_len < 48)
     {
-      pad (&w[ 8], pw_len & 0xf);
+      pad (&w[ 8], out_len & 0xf);
     }
-    else if (pw_len < 64)
+    else if (out_len < 64)
     {
-      pad (&w[12], pw_len & 0xf);
+      pad (&w[12], out_len & 0xf);
     }
 
-    domino_big_md (w, pw_len, state, s_lotus_magic_table);
+    domino_big_md (w, out_len, state, s_lotus_magic_table);
 
     const u32 w0_t = uint_to_hex_upper8 ((state[0] >>  0) & 255) <<  0
                      | uint_to_hex_upper8 ((state[0] >>  8) & 255) << 16;
@@ -626,24 +626,24 @@ __kernel void m08700_s04 (__global pw_t *pws, __global kernel_rule_t *rules_buf,
      * padding
      */
 
-    if (pw_len < 16)
+    if (out_len < 16)
     {
-      pad (&w[ 0], pw_len & 0xf);
+      pad (&w[ 0], out_len & 0xf);
     }
-    else if (pw_len < 32)
+    else if (out_len < 32)
     {
-      pad (&w[ 4], pw_len & 0xf);
+      pad (&w[ 4], out_len & 0xf);
     }
-    else if (pw_len < 48)
+    else if (out_len < 48)
     {
-      pad (&w[ 8], pw_len & 0xf);
+      pad (&w[ 8], out_len & 0xf);
     }
-    else if (pw_len < 64)
+    else if (out_len < 64)
     {
-      pad (&w[12], pw_len & 0xf);
+      pad (&w[12], out_len & 0xf);
     }
 
-    domino_big_md (w, pw_len, state, s_lotus_magic_table);
+    domino_big_md (w, out_len, state, s_lotus_magic_table);
 
     const u32 w0_t = uint_to_hex_upper8 ((state[0] >>  0) & 255) <<  0
                      | uint_to_hex_upper8 ((state[0] >>  8) & 255) << 16;