From b6f10a2a81796ed29dd83a8ccb4b334600c9e7cc Mon Sep 17 00:00:00 2001 From: Jens Steube Date: Thu, 4 Feb 2016 23:25:52 +0100 Subject: [PATCH] Fix some bugs related to -a0, all caused by the same pitfall --- OpenCL/m04700_a0.cl | 8 ++++++-- OpenCL/m06900_a0.cl | 4 ++-- OpenCL/m07600_a0.cl | 4 ++-- OpenCL/m08300_a0.cl | 4 ++-- OpenCL/m08400_a0.cl | 4 ++-- OpenCL/m08600_a0.cl | 36 ++++++++++++++++++------------------ OpenCL/m08700_a0.cl | 36 ++++++++++++++++++------------------ 7 files changed, 50 insertions(+), 46 deletions(-) diff --git a/OpenCL/m04700_a0.cl b/OpenCL/m04700_a0.cl index 1854e1944..8a9b5a291 100644 --- 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 a0f13d297..1b1563224 100644 --- 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 de87dc681..5862cdae2 100644 --- 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 2872e9f1f..3346c6489 100644 --- 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 348c994e3..2c29b4138 100644 --- 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 1c688fad8..bae3f4c86 100644 --- 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 63764228b..91785fcc9 100644 --- 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;